/opencl/tools

#include <stdio.h>
#include <stdlib.h>
#include <string.h>
#include <pthread.h>

#include <CL/cl.h>

#define ITERATIONS 500000

#define fail(code, reason) {\
    fprintf(stderr, "%s\n", reason); \
    exit (code); \
}

pthread_barrier_t bar;

typedef struct  {
    cl_uint num_devices;
    cl_device_id *devices;
} device_info_t;

static inline cl_program load_app(cl_context ctx, cl_uint num_devices, cl_device_id *devices, const char *name) {
    FILE *f;
    size_t len;
    char *buf = "__kernel void test() {}";
    device_info_t info = { num_devices, devices }; 
    
    cl_program app;
    
    if (name) {
	f = fopen(name, "rb");
	if (!f) return NULL;

        fseek(f, 0, SEEK_END); 
	len = ftell(f);
	fseek(f, 0, SEEK_SET); 
    
	buf = (char*)malloc(len*sizeof(char));
	if (!buf) return NULL;
    
	if (fread(buf, 1, len, f) != len) {
	    free(buf);
	    fclose(f);
	    return NULL;
	}
    
	fclose(f);
    } 

    int i;
    printf("Compiling for ");
    for (i=0;i<num_devices;i++) {
	char name[256];
	int err = clGetDeviceInfo(devices[i], CL_DEVICE_NAME, sizeof(name), name, NULL);
	if (err != CL_SUCCESS) puts("Unknown ");
	else printf("%s ", name);
    }
    printf("\n");
    
    app = clCreateProgramWithSource(ctx, 1, (const char**)&buf, &len, NULL);
    
    
    if (app) {
	size_t size;
	char build_log[4096];
	cl_build_status build_status;

	//"-cl-nv-maxrregcount=48"
	//"-cl-nv-opt-level=<N>" (0 - no optimizations, 3 - default)
	//"-cl-nv-arch sm_XX" - selects the target CUDA Compute Level architecture to compile for (sm_10 for 1.0, sm_11 for 1.1, sm_12 for 1.2, sm_13 for 1.3 and sm_20 for 2.0 (Fermi))
	//"--cl-nv-cstd=CLX.X" - selects the target OpenCL C language version (CL1.0 or CL1.1)
	char *build_flags = "-cl-nv-arch sm_20";
	int err = clBuildProgram(app, num_devices, devices, build_flags, NULL, NULL);
	if (err != CL_SUCCESS) printf("Application build failed (%i)\n", err);
	    
	for (i = 0; i < num_devices; i++) {
	    do {
		err = clGetProgramBuildInfo(app, devices[i], CL_PROGRAM_BUILD_STATUS, sizeof(cl_build_status), &build_status, NULL);
	    } while (build_status == CL_BUILD_IN_PROGRESS);
		
	    err = clGetProgramBuildInfo(app, devices[i], CL_PROGRAM_BUILD_LOG, sizeof(build_log) - 1, &build_log, &size);
	    if (!err) {
		build_log[size] = 0;
	        if (!err) {
		    if (build_status == CL_BUILD_SUCCESS) {
		    } else if (build_status == CL_BUILD_ERROR) {
		        printf("Build failed for device %i:\n======================\n%s\n--------------------\n\n", i, build_log);
		    } else {
		        printf("Build failed for device %i\n", i);
		    }
		}
	    }
	}
    } else {
	printf("Program creation failed\n");
    }
        
    if (name) free(buf);
    return app;
}

void bench(cl_context ctx, cl_device_id device, cl_kernel kern) {
    cl_int err;

    size_t runtime;
    size_t us, flops;
    struct timeval tv1,tv2;
    size_t i, iterations = ITERATIONS;
    size_t dims[2] = {1024, 1024};

    cl_command_queue queue = clCreateCommandQueue(ctx, device, 0, &err);
    if (err != CL_SUCCESS) 
	fail(-1, "clCreateCommandQueue() failed");

    clEnqueueNDRangeKernel (queue, kern, 2, NULL, dims, NULL, 0, NULL, NULL);
    clFinish(queue);

    // synchronizing
    pthread_barrier_wait(&bar);

    gettimeofday(&tv1, NULL);
    for (i = 0; i < iterations; i++) {
	clEnqueueNDRangeKernel (queue, kern, 2, NULL, dims, NULL, 0, NULL, NULL);
    }
    clFinish(queue);
    gettimeofday(&tv2, NULL);

    us = (tv2.tv_sec - tv1.tv_sec)*1000000 + (tv2.tv_usec - tv1.tv_usec);
    us /= iterations;


    clReleaseCommandQueue(queue);

    printf("  Execution time: %lu us\n", us);

}

struct thread_info {
    cl_context ctx;
    cl_device_id device;
    cl_kernel kern;
};

void *bench_thread(void *vargs) {
    struct thread_info *args = (struct thread_info*)vargs;

    bench(args->ctx, args->device, args->kern);
}


int main(int argc, char *argv[]) {
    int devid = -1;
    char *source = NULL;
    
    cl_int err;
    cl_device_id devices[16];
    pthread_t thr[16];
    struct thread_info args[16];

    cl_context ctx;
    cl_context ctxs[16];

    cl_program app;
    cl_program apps[16];

    cl_kernel kern;
    cl_kernel kerns[16];


    char fname[256];    
    unsigned char *binary[16];
    size_t binary_size[16];
    size_t real_size;


    FILE *f;
    
/*
    if (argc < 2) {
	fprintf(stderr, "Usage: %s [source] [device]\n", argv[0]);
	exit(0);	
    }
    

    if (argc > 2) {
	devid = atoi(argv[2]);	
    }
*/

    cl_uint max = 0;
    cl_platform_id selected_platform;
    cl_uint i, j, num_platforms, num_devices;
    cl_platform_id platforms[4];

    clGetPlatformIDs(4, platforms, &num_platforms);

    for (i = 0; i < num_platforms; i++) {
	err = clGetDeviceIDs(platforms[i], CL_DEVICE_TYPE_GPU, 0, NULL, &num_devices);
	if ((err == CL_SUCCESS)&&(num_devices > max)) {
	    max = num_devices;
	    selected_platform = platforms[i];
	}
    }
    
    clGetDeviceIDs(selected_platform, CL_DEVICE_TYPE_GPU, 16, devices, &num_devices);
    
    if (devid >= (int)num_devices) fail(-1, "Invalid device number is specified");


//    cl_context ctx = clCreateContextFromType(NULL, CL_DEVICE_TYPE_GPU, NULL, NULL, NULL);
//    if (!ctx) fail(-1, "Failed to create OpenCL context");

  
//    err = clGetContextInfo(ctx, CL_CONTEXT_DEVICES, 16 * sizeof(cl_device_id), &devices, &num_devices);
//    if (err != CL_SUCCESS) fail(-1, "clGetContextInfo call is failed");

//    num_devices /= sizeof(cl_device_id);

    
    for (i = 0; i < num_devices; i++) {
	ctxs[i] = clCreateContext(0, 1, &devices[i], NULL, NULL, &err);
        if (!ctxs[i]) fail(-1, "Failed to create OpenCL context");
	
	apps[i] = load_app(ctxs[i], 1, &devices[i], source);
	if (!apps[i]) fail(-1, "Compilation failed");
    
	kerns[i] = clCreateKernel(apps[i], "test", &err);
	if (err != CL_SUCCESS) fail(-1, "Error creating kernel");
    }

    pthread_barrier_init(&bar, NULL, num_devices);

    printf("Running in individual context:\n");
    for (i = 0; i < num_devices; i++) {
	args[i].ctx = ctxs[i];
	args[i].device = devices[i];
	args[i].kern = kerns[i];
//	pthread_create(&thr[i], NULL, bench_thread, &args[i]);
    }

    for (i = 0; i < num_devices; i++) {
	void *res;
//	pthread_join(thr[i], &res);
        clReleaseKernel(kerns[i]);
	clReleaseProgram(apps[i]);
	clReleaseContext(ctxs[i]);
    }
    printf("\n\n");

    pthread_barrier_destroy(&bar);

    for (j = /*2*/num_devices; j <= num_devices; j++) {
	pthread_barrier_init(&bar, NULL, j);
	
	ctx = clCreateContext(0, j, devices, NULL, NULL, &err);
	if (!ctx) fail(-1, "Failed to create OpenCL context");

	app = load_app(ctx, j, devices, source);
	if (!app) fail(-1, "Compilation failed");

	kern = clCreateKernel(app, "test", &err);
	if (err != CL_SUCCESS) fail(-1, "Error creating kernel");


	printf("Running parallel for %i devices:\n", j);
	for (i = 0; i < j; i++) {
	    args[i].ctx = ctx;
	    args[i].device = devices[i];
	    args[i].kern = kern;
	    pthread_create(&thr[i], NULL, bench_thread, &args[i]);
	}

	for (i = 0; i < j; i++) {
	    void *res;
	    pthread_join(thr[i], &res);
	}

	printf("\n\n");

	pthread_barrier_destroy(&bar);
    
	clReleaseKernel(kern);
	clReleaseProgram(app);
	clReleaseContext(ctx);
    }
    
    pthread_barrier_init(&bar, NULL, 1);

    ctx = clCreateContext(0, num_devices, devices, NULL, NULL, &err);
    if (!ctx) fail(-1, "Failed to create OpenCL context");

    app = load_app(ctx, num_devices, devices, source);
    if (!app) fail(-1, "Compilation failed");

    kern = clCreateKernel(app, "test", &err);
    if (err != CL_SUCCESS) fail(-1, "Error creating kernel");

    printf("Running sequential (multi-devices context):\n");
    for (i = 0; i < num_devices; i++) {
	bench(ctx, devices[i], kern);
    }

    clReleaseKernel(kern);
    clReleaseProgram(app);
    clReleaseContext(ctx);
    
    pthread_barrier_destroy(&bar);
}