/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);

}