sampler* initialize_sampler(cl_int chain_length, cl_int dimension,
cl_int walkers_per_group, size_t work_group_size,
double a, cl_int pdf_number,
cl_int data_length, cl_float *data,
cl_int num_to_save, cl_int *indices_to_save,
const char *plat_name, const char *dev_name){
/*
Initialize stretch move MCMC sampler struct.
Arrange parameters into sampler struct pointer.
Allocate arrays on host, initialize walkers and other values as appropriate.
Start OpenCL context and queue.
Allocate device memory and transfer from host.
Compile and initialize random number generator.
Compile stretch move OpenCL kernel.
Input:
cl_int chain_length Allocate space for this many samples in the sampler struct.
Sampler fills this array when run_sampler is called.
cl_int dimension Dimension of state vector of Markov chain.
cl_int walkers_per_group Number of walkers in each of two groups. Total walkers is twice this.
size_t work_group_size Work group size.
For CPU this must be set to one.
For GPU this should be set larger, powers of two are optimal, try 64, 128 or 256.
This number must divide walkers_per_group.
double a Coefficient for range of 'z' random variable.
Must be greater than one.
Standard value is 2.
Decrease a to increase low acceptance rate, especially in high dimensions.
cl_int pdf_number Which PDF to sample. Passed to pdf.h as a compile time definition.
cl_int data_length Length of observation data. If no data set this to zero.
cl_float *data Observation data.
cl_int num_to_save Number of components to save in the chain
cl_int *indices_to_save Indices of components to save in the chain
const char *plat_name String for platform name. Set to CHOOSE_INTERACTIVELY (no quotes) to do so.
const char *dev_name String for device name. Set to CHOOSE_INTERACTIVELY (no quotes) to do so.
Output:
returned: sampler *samp Pointer to sampler struct with parameters, arrays, context, queue, kernel initialized.
*/
if(OUTPUT_LEVEL > 0) printf("Initializing Stretch Move sampler.\n");
// --------------------------------------------------------------------------
// Set parameters
// --------------------------------------------------------------------------
// This environment variable forces headers to be reloaded each time
// If not set and pdf if changed, changes may not be updated
setenv("CUDA_CACHE_DISABLE", "1", 1);
// allocate the structure for all the sampler parameters and arrays
sampler * samp = (sampler *) malloc(sizeof(sampler));
if(!samp) { perror("Allocation failure sampler"); abort(); }
// user set parameters
samp->M = chain_length; // Number of steps to run
samp->N = dimension; // Dimension of the problem and the walkers
samp->K_over_two = walkers_per_group ; // Number of walkers in each group
// derived parameters
samp->K = 2 * samp->K_over_two; // Total walkers
samp->total_samples = samp->M * samp->K; // Total samples produced
// indices to save
samp->num_to_save = num_to_save;
samp->indices_to_save_host = indices_to_save;
// Allocate the structure and set values
samp->data_st = (data_struct *) malloc(sizeof(data_struct));
if(!(samp->data_st)) { perror("Allocation failure data_struct"); abort(); }
// default value one, unless performing simulated annealing
(samp->data_st)->beta = 1.0f;
(samp->data_st)->save = 1;
(samp->data_st)->num_to_save = num_to_save;
// coefficient on Z random variable
samp->a = a;
double a_coeffs[3];
a_coeffs[0] = 1.0 / a;
a_coeffs[1] = 2.0 * (1.0 - 1.0/a);
a_coeffs[2] = a - 2.0 + 1.0/a;
// error check on dimensions
if(samp->K <= samp->N){
fprintf(stderr, "Error: Must have more walkers than the dimension.\nExiting\n");
abort();
}
// error check on work sizes
if( (samp->K_over_two % work_group_size) != 0){
fprintf(stderr, "Error: Number of walkers in each group must be multiple of work group size.\nExiting\n");
abort();
}
// error check on dimensions to save
for(int i=0; i<num_to_save; i++){
if(samp->indices_to_save_host[i] >= samp->N){
fprintf(stderr, "Error: Cannot save an index larger than the dimension of the problem.\nExiting\n");
abort();
}
}
if(a <= 1.0){
fprintf(stderr, "Error: Value of a must be greater than one.\nDefaulting to 2.\n");
samp->a = 2.0;
}
// for later output
samp->acor_times = (double *) malloc(samp->num_to_save * sizeof(double));
if(!samp->acor_times) { perror("Allocation failure"); abort(); }
samp->acor_pass = (char *) malloc(samp->num_to_save * sizeof(char));
if(!samp->acor_pass) { perror("Allocation failure"); abort(); }
samp->sigma = (double *) malloc(samp->num_to_save * sizeof(double));
if(!samp->sigma) { perror("Allocation failure"); abort(); }
samp->means = (double *) malloc(samp->num_to_save * sizeof(double));
if(!samp->means) { perror("Allocation failure"); abort(); }
samp->err_bar = (double *) malloc(samp->num_to_save * sizeof(double));
if(!samp->err_bar) { perror("Allocation failure"); abort(); }
// write parameter file for plotting
write_parameter_file_matlab(samp->M, samp->N, samp->K, "Stretch Move",
samp->indices_to_save_host, samp->num_to_save, pdf_number);
// --------------------------------------------------------------------------
// Set up OpenCL context and queues
// --------------------------------------------------------------------------
if(OUTPUT_LEVEL > 0) printf("Begin opencl contexts.\n");
create_context_on(plat_name, dev_name, 0, &(samp->ctx), NULL, 0);
{
cl_int status;
cl_device_id my_dev;
CALL_CL_GUARDED(clGetContextInfo, (samp->ctx, CL_CONTEXT_DEVICES,
sizeof(my_dev), &my_dev, NULL));
samp->queue = clCreateCommandQueue(samp->ctx, my_dev, 0, &status);
CHECK_CL_ERROR(status, "clCreateCommandQueue");
samp->queue_mem = clCreateCommandQueue(samp->ctx, my_dev, 0, &status);
CHECK_CL_ERROR(status, "clCreateCommandQueue");
}
// print information on selected device
if(OUTPUT_LEVEL > 1) print_device_info_from_queue(samp->queue);
// set the work group sizes
samp->ldim[0] = work_group_size;
samp->gdim[0] = samp->K_over_two;
if(OUTPUT_LEVEL > 0) printf("Context built.\n");
// --------------------------------------------------------------------------
// Start total timing
// --------------------------------------------------------------------------
if(OUTPUT_LEVEL > 0) printf("Begin total timing.\n");
get_timestamp(&(samp->time1_total));
// --------------------------------------------------------------------------
// Allocate host memory
// --------------------------------------------------------------------------
// counter for number of samples accepted
samp->accepted_host = (cl_ulong *) malloc(samp->K_over_two * sizeof(cl_ulong));
if(!(samp->accepted_host)){ perror("Allocation failure accepted host"); abort(); }
for(int i=0; i< (samp->K_over_two); i++) samp->accepted_host[i] = 0;
// Adjacent memory on x_red moves with in the walker
// To access the ith component of walker j, take x_red[i + j*N];
// red walkers
samp->X_red_host = (cl_float *) malloc(samp->N * samp->K_over_two * sizeof(cl_float));
if(!(samp->X_red_host)){ perror("Allocation failure X_red_host"); abort(); }
// log likelihood
samp->log_pdf_red_host = (cl_float *) malloc(samp->K_over_two * sizeof(cl_float));
if(!(samp->log_pdf_red_host)){ perror("Allocation failure X_red_host"); abort(); }
for(int i=0; i<(samp->K_over_two); i++) samp->log_pdf_red_host[i] = (-1.0f) / 0.0f;
// black walkers
samp->X_black_host = (cl_float *) malloc(samp->N * samp->K_over_two * sizeof(cl_float));
if(!(samp->X_black_host)){ perror("Allocation failure X_black_host"); abort(); }
// log likelihood
samp->log_pdf_black_host = (cl_float *) malloc(samp->K_over_two * sizeof(cl_float));
if(!(samp->log_pdf_black_host)){ perror("Allocation failure X_red_host"); abort(); }
for(int i=0; i< (samp->K_over_two); i++) samp->log_pdf_black_host[i] = (-1.0f) / 0.0f;
// samples on host
cl_int samples_length = samp->num_to_save * samp->M * samp->K; // length of the samples array
samp->samples_host = (cl_float *) malloc(samples_length * sizeof(cl_float)); // samples to return
if(!(samp->samples_host)){ perror("Allocation failure samples_host"); abort(); }
// intialize the walkers to random values
// set the seed value
srand48(0);
// initialize the walkers to small random values
for(int j=0; j < samp->N * samp->K_over_two; j++){
if(NONNEGATIVE_BOX){
samp->X_black_host[j] = (cl_float) drand48();
samp->X_red_host[j] = (cl_float) drand48();
}
else{
samp->X_black_host[j] = (cl_float) (0.1 * (drand48()-0.5));
samp->X_red_host[j] = (cl_float) (0.1 * (drand48()-0.5));
}
}
// set up observations
samp->data_length = data_length;
// there are lots of complications that appear if this is empty
// make it length one instead
if(samp->data_length == 0){
samp->data_length = 1;
samp->data_host = (cl_float *) malloc(samp->data_length * sizeof(cl_float)) ;
if(!(samp->data_host)){ perror("Allocation failure data_host"); abort(); }
samp->data_host[0] = 0.0f;
}
else{
// standard case
samp->data_host = data;
}
// --------------------------------------------------------------------------
// load kernels
// --------------------------------------------------------------------------
// stretch move kernel
char *knl_text = read_file("stretch_move.cl");
char options[300];
sprintf(options, "-D NN=%d -D K_OVER_TWO=%d -D WORK_GROUP_SIZE=%d -D DATA_LEN=%d -D PDF_NUMBER=%d -D A_COEFF_0=%.10ff -D A_COEFF_1=%.10ff -D A_COEFF_2=%.10ff -I . ",
samp->N, samp->K_over_two, (int) work_group_size, samp->data_length, pdf_number, a_coeffs[0], a_coeffs[1], a_coeffs[2]);
if(OUTPUT_LEVEL > 0) printf("Options string for stretch move kernel:%s\n", options);
samp->stretch_knl = kernel_from_string(samp->ctx, knl_text, "stretch_move", options);
free(knl_text);
if(OUTPUT_LEVEL > 0) printf("Stretch Move kernel compiled.\n");
// random number generator initialization
char * knl_text_rand = read_file("Kernel_Ranluxcl_Init.cl");
char options_rand_lux[100];
if(AMD)
sprintf(options_rand_lux, "-DRANLUXCL_LUX=4 -I .");
else
sprintf(options_rand_lux, "-DRANLUXCL_LUX=4");
samp->init_rand_lux_knl = kernel_from_string(samp->ctx, knl_text_rand, "Kernel_Ranluxcl_Init", options_rand_lux);
free(knl_text_rand);
if(OUTPUT_LEVEL > 0) printf("Ranluxcl init kernel compiled.\n");
// --------------------------------------------------------------------------
// allocate device memory
// --------------------------------------------------------------------------
cl_int status;
samp->X_red_device = clCreateBuffer(samp->ctx, CL_MEM_READ_WRITE,
sizeof(cl_float) * samp->N * samp->K_over_two, 0, &status);
CHECK_CL_ERROR(status, "clCreateBuffer");
samp->log_pdf_red_device = clCreateBuffer(samp->ctx, CL_MEM_READ_WRITE,
sizeof(cl_float) * samp->K_over_two, 0, &status);
CHECK_CL_ERROR(status, "clCreateBuffer");
samp->X_red_save = clCreateBuffer(samp->ctx, CL_MEM_WRITE_ONLY,
sizeof(cl_float) * samp->num_to_save * samp->K_over_two, 0, &status);
CHECK_CL_ERROR(status, "clCreateBuffer");
samp->X_black_device = clCreateBuffer(samp->ctx, CL_MEM_READ_WRITE,
sizeof(cl_float) * samp->N * samp->K_over_two, 0, &status);
CHECK_CL_ERROR(status, "clCreateBuffer");
samp->log_pdf_black_device = clCreateBuffer(samp->ctx, CL_MEM_READ_WRITE,
sizeof(cl_float) * samp->K_over_two, 0, &status);
CHECK_CL_ERROR(status, "clCreateBuffer");
samp->X_black_save = clCreateBuffer(samp->ctx, CL_MEM_WRITE_ONLY,
sizeof(cl_float) * samp->num_to_save * samp->K_over_two, 0, &status);
CHECK_CL_ERROR(status, "clCreateBuffer");
samp->accepted_device = clCreateBuffer(samp->ctx, CL_MEM_READ_WRITE,
samp->K_over_two * sizeof(cl_ulong), 0, &status);
CHECK_CL_ERROR(status, "clCreateBuffer");
samp->indices_to_save_device = clCreateBuffer(samp->ctx, CL_MEM_READ_ONLY,
samp->num_to_save * sizeof(cl_int), 0, &status);
CHECK_CL_ERROR(status, "clCreateBuffer");
// allocate for the observations
samp->data_device = clCreateBuffer(samp->ctx, CL_MEM_READ_WRITE,
sizeof(cl_float) * samp->data_length, 0, &status);
CHECK_CL_ERROR(status, "clCreateBuffer");
// data struct on device
samp->data_st_device = clCreateBuffer(samp->ctx, CL_MEM_READ_WRITE,
sizeof(data_struct), 0, &status);
CHECK_CL_ERROR(status, "clCreateBuffer");
// allocate for the state array for randluxcl
// use a 1d work group
size_t rand_lux_state_buffer_size = samp->gdim[0] * 7 * sizeof(cl_float4);
samp->ranluxcltab = clCreateBuffer(samp->ctx, CL_MEM_READ_WRITE,
rand_lux_state_buffer_size, 0, &status);
CHECK_CL_ERROR(status, "clCreateBuffer");
// --------------------------------------------------------------------------
// transfer to device
// --------------------------------------------------------------------------
CALL_CL_GUARDED(clEnqueueWriteBuffer, (
samp->queue, samp->X_red_device, /*blocking*/ CL_TRUE, /*offset*/ 0,
samp->N * samp->K_over_two * sizeof(cl_float), samp->X_red_host,
0, NULL, NULL));
CALL_CL_GUARDED(clEnqueueWriteBuffer, (
samp->queue, samp->log_pdf_red_device, /*blocking*/ CL_TRUE, /*offset*/ 0,
samp->K_over_two * sizeof(cl_float), samp->log_pdf_red_host,
0, NULL, NULL));
CALL_CL_GUARDED(clEnqueueWriteBuffer, (
samp->queue, samp->X_black_device, /*blocking*/ CL_TRUE, /*offset*/ 0,
samp->N * samp->K_over_two * sizeof(cl_float), samp->X_black_host,
0, NULL, NULL));
CALL_CL_GUARDED(clEnqueueWriteBuffer, (
samp->queue, samp->log_pdf_black_device, /*blocking*/ CL_TRUE, /*offset*/ 0,
samp->K_over_two * sizeof(cl_float), samp->log_pdf_black_host,
0, NULL, NULL));
CALL_CL_GUARDED(clEnqueueWriteBuffer, (
samp->queue, samp->data_device, /*blocking*/ CL_TRUE, /*offset*/ 0,
samp->data_length * sizeof(cl_float), samp->data_host,
0, NULL, NULL));
CALL_CL_GUARDED(clEnqueueWriteBuffer, (
samp->queue, samp->data_st_device, /*blocking*/ CL_TRUE, /*offset*/ 0,
sizeof(data_struct), samp->data_st,
0, NULL, NULL));
CALL_CL_GUARDED(clEnqueueWriteBuffer, (
samp->queue, samp->indices_to_save_device, /*blocking*/ CL_TRUE, /*offset*/ 0,
samp->num_to_save * sizeof(cl_int), samp->indices_to_save_host,
0, NULL, NULL));
CALL_CL_GUARDED(clFinish, (samp->queue));
// --------------------------------------------------------------------------
// Initialize random number generator
// --------------------------------------------------------------------------
// int for state variable initialization
cl_int ins = 1;
SET_2_KERNEL_ARGS(samp->init_rand_lux_knl, ins, samp->ranluxcltab);
CALL_CL_GUARDED(clEnqueueNDRangeKernel,
(samp->queue, samp->init_rand_lux_knl,
/*dimensions*/ 1, NULL, samp->gdim, samp->ldim,
0, NULL, NULL));
CALL_CL_GUARDED(clFinish, (samp->queue));
// --------------------------------------------------------------------------
// Sampler initialization is done
// --------------------------------------------------------------------------
if(OUTPUT_LEVEL > 0) printf("Sampler initialized.\n");
return samp;
}
int main(int argc, char **argv)
{
if (argc != 3)
{
fprintf(stderr, "need two arguments!\n");
abort();
}
const long n = atol(argv[1]);
const long size = n*n;
const int ntrips = atoi(argv[2]);
cl_context ctx;
cl_command_queue queue;
create_context_on(CHOOSE_INTERACTIVELY, CHOOSE_INTERACTIVELY, 0, &ctx, &queue, 0);
cl_int status;
// --------------------------------------------------------------------------
// load kernels
// --------------------------------------------------------------------------
char *knl_text = read_file("transpose-soln.cl");
cl_kernel knl = kernel_from_string(ctx, knl_text, "transpose", NULL);
free(knl_text);
// --------------------------------------------------------------------------
// allocate and initialize CPU memory
// --------------------------------------------------------------------------
#ifdef USE_PINNED
cl_mem buf_a_host = clCreateBuffer(ctx,
CL_MEM_READ_WRITE | CL_MEM_ALLOC_HOST_PTR,
sizeof(value_type) * size, 0, &status);
CHECK_CL_ERROR(status, "clCreateBuffer");
cl_mem buf_b_host = clCreateBuffer(ctx,
CL_MEM_READ_WRITE | CL_MEM_ALLOC_HOST_PTR,
sizeof(value_type) * size, 0, &status);
CHECK_CL_ERROR(status, "clCreateBuffer");
value_type *a = (value_type *) clEnqueueMapBuffer(queue, buf_a_host,
/*blocking*/ CL_TRUE, CL_MAP_WRITE_INVALIDATE_REGION,
/*offs*/ 0, sizeof(value_type)*size, 0, NULL, NULL, &status);
CHECK_CL_ERROR(status, "clEnqueueMapBuffer");
value_type *b = (value_type *) clEnqueueMapBuffer(queue, buf_b_host,
/*blocking*/ CL_TRUE, CL_MAP_WRITE_INVALIDATE_REGION,
/*offs*/ 0, sizeof(value_type)*size, 0, NULL, NULL, &status);
CHECK_CL_ERROR(status, "clEnqueueMapBuffer");
#else
value_type *a = (value_type *) malloc(sizeof(value_type) * size);
if (!a) { perror("alloc x"); abort(); }
value_type *b = (value_type *) malloc(sizeof(value_type) * size);
if (!b) { perror("alloc y"); abort(); }
#endif
for (size_t j = 0; j < n; ++j)
for (size_t i = 0; i < n; ++i)
a[i + j*n] = i + j*n;
// --------------------------------------------------------------------------
// allocate device memory
// --------------------------------------------------------------------------
cl_mem buf_a = clCreateBuffer(ctx, CL_MEM_READ_WRITE,
sizeof(value_type) * size, 0, &status);
CHECK_CL_ERROR(status, "clCreateBuffer");
cl_mem buf_b = clCreateBuffer(ctx, CL_MEM_READ_WRITE,
sizeof(value_type) * size, 0, &status);
CHECK_CL_ERROR(status, "clCreateBuffer");
// --------------------------------------------------------------------------
// transfer to device
// --------------------------------------------------------------------------
CALL_CL_GUARDED(clFinish, (queue));
timestamp_type time1, time2;
get_timestamp(&time1);
CALL_CL_GUARDED(clEnqueueWriteBuffer, (
queue, buf_a, /*blocking*/ CL_FALSE, /*offset*/ 0,
size * sizeof(value_type), a,
0, NULL, NULL));
CALL_CL_GUARDED(clEnqueueWriteBuffer, (
queue, buf_b, /*blocking*/ CL_FALSE, /*offset*/ 0,
size * sizeof(value_type), b,
0, NULL, NULL));
get_timestamp(&time2);
double elapsed = timestamp_diff_in_seconds(time1,time2);
printf("transfer: %f s\n", elapsed);
printf("transfer: %f GB/s\n",
2*size*sizeof(value_type)/1e9/elapsed);
// --------------------------------------------------------------------------
// run code on device
// --------------------------------------------------------------------------
CALL_CL_GUARDED(clFinish, (queue));
get_timestamp(&time1);
for (int trip = 0; trip < ntrips; ++trip)
{
SET_3_KERNEL_ARGS(knl, buf_a, buf_b, n);
size_t ldim[] = { 16, 16 };
size_t gdim[] = { n, n };
CALL_CL_GUARDED(clEnqueueNDRangeKernel,
(queue, knl,
/*dimensions*/ 2, NULL, gdim, ldim,
0, NULL, NULL));
}
CALL_CL_GUARDED(clFinish, (queue));
get_timestamp(&time2);
elapsed = timestamp_diff_in_seconds(time1,time2)/ntrips;
printf("%f s\n", elapsed);
printf("%f GB/s\n",
2*size*sizeof(value_type)/1e9/elapsed);
CALL_CL_GUARDED(clEnqueueReadBuffer, (
queue, buf_b, /*blocking*/ CL_FALSE, /*offset*/ 0,
size * sizeof(value_type), b,
0, NULL, NULL));
CALL_CL_GUARDED(clFinish, (queue));
for (size_t i = 0; i < n; ++i)
for (size_t j = 0; j < n; ++j)
if (a[i + j*n] != b[j + i*n])
{
printf("bad %d %d\n", i, j);
abort();
}
// --------------------------------------------------------------------------
// clean up
// --------------------------------------------------------------------------
CALL_CL_GUARDED(clFinish, (queue));
CALL_CL_GUARDED(clReleaseMemObject, (buf_a));
CALL_CL_GUARDED(clReleaseMemObject, (buf_b));
CALL_CL_GUARDED(clReleaseKernel, (knl));
CALL_CL_GUARDED(clReleaseCommandQueue, (queue));
CALL_CL_GUARDED(clReleaseContext, (ctx));
#ifdef USE_PINNED
CALL_CL_GUARDED(clReleaseMemObject, (buf_a_host));
CALL_CL_GUARDED(clReleaseMemObject, (buf_b_host));
#else
free(a);
free(b);
#endif
return 0;
}
int
main (int argv, char **argc)
{
/////////////////////////
////// SAME IN EVERY FILE
/////////////////////////
// create context and command queue
cl_context __sheets_context;
cl_command_queue __sheets_queue;
int _i;
cl_int __cl_err;
create_context_on(SHEETS_PLAT_NAME,
SHEETS_DEV_NAME,
0, /* choose the first (only) available device */
&__sheets_context,
&__sheets_queue,
0);
// compile kernels
for (_i = 0; _i < NKERNELS; _i++) {
compiled_kernels[_i] = kernel_from_string(__sheets_context,
kernel_strings[_i],
kernel_names[_i],
SHEETS_KERNEL_COMPILE_OPTS);
}
////// [END]
size_t __SIZE_wav = atoi(argc[1]);
float wav[__SIZE_wav];
const char *file_name = "mytune.wav";
int in_thrsh_cnt = 0;
timestamp_type st;
timestamp_type end;
get_timestamp(&st);
for (_i = 0; _i < __SIZE_wav; _i++) {
wav[_i] = (float) rand() / RAND_MAX;
if (in_thrsh(wav[_i], 0.1112, 0.7888))
in_thrsh_cnt++;
}
get_timestamp(&end);
printf("cpu execution took %f seconds\n", timestamp_diff_in_seconds(st, end));
get_timestamp(&st);
/////////////////
////// GFUNC CALL
/////////////////
/// create variables for function arguments given as literals
float __PRIM_band_restrict_ARG2 = 0.1112f;
float __PRIM_band_restrict_ARG3 = 0.7888f;
/// return array (always arg0)
cl_mem __CLMEM_band_restrict_ARG0 = clCreateBuffer(__sheets_context,
CL_MEM_WRITE_ONLY,
sizeof(float) * __SIZE_wav,
NULL,
&__cl_err);
CHECK_CL_ERROR(__cl_err, "clCreateBuffer");
/// input arrays
cl_mem __CLMEM_band_restrict_ARG1 = clCreateBuffer(__sheets_context,
CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
sizeof(float) * __SIZE_wav,
(void *) wav,
&__cl_err);
CHECK_CL_ERROR(__cl_err, "clCreateBuffer");
/// write to device memory
CALL_CL_GUARDED(clEnqueueWriteBuffer,
(__sheets_queue,
__CLMEM_band_restrict_ARG1,
CL_TRUE, /* blocking write */
0, /* no offset */
sizeof(float) * __SIZE_wav,
wav,
0, /* no wait list */
NULL,
NULL)
);
/// set up kernel arguments
SET_4_KERNEL_ARGS(compiled_kernels[0],
__CLMEM_band_restrict_ARG0,
__CLMEM_band_restrict_ARG1,
__PRIM_band_restrict_ARG2,
__PRIM_band_restrict_ARG3);
/// enqueue kernel
cl_event __CLEVENT_band_restrict_CALL;
CALL_CL_GUARDED(clEnqueueNDRangeKernel,
(__sheets_queue,
compiled_kernels[0],
1, /* 1 dimension */
0, /* 0 offset */
&__SIZE_wav,
NULL, /* let OpenCL break things up */
0, /* no events in wait list */
NULL, /* empty wait list */
&__CLEVENT_band_restrict_CALL)
);
/// allocate space for cpu return array
float out[__SIZE_wav];
CALL_CL_GUARDED(clEnqueueReadBuffer,
(__sheets_queue,
__CLMEM_band_restrict_ARG0,
CL_TRUE, /* blocking read */
0, /* 0 offset */
sizeof(float) * __SIZE_wav, /* read whole buffer */
(void *) out, /* host pointer */
1, /* wait for gfunc to finish */
&__CLEVENT_band_restrict_CALL, /* "" */
NULL) /* no need to wait for this call though */
);
////// [END] GFUNC CALL
get_timestamp(&end);
printf("gfunc call took %f seconds\n", timestamp_diff_in_seconds(st, end));
////// Validate call
int c = 0;
for (_i = 0; _i < __SIZE_wav; _i++) {
if (in_thrsh(out[_i], 0.1112, 0.7888)) {
c++;
} else if(out[_i]) {
exit(1);
}
}
printf("\n");
assert(in_thrsh_cnt == c);
//////////////
////// CLEANUP
//////////////
CALL_CL_GUARDED(clReleaseMemObject, (__CLMEM_band_restrict_ARG0));
CALL_CL_GUARDED(clReleaseMemObject, (__CLMEM_band_restrict_ARG1));
for (_i = 0; _i < NKERNELS; _i++) {
CALL_CL_GUARDED(clReleaseKernel, (compiled_kernels[_i]));
}
CALL_CL_GUARDED(clReleaseCommandQueue, (__sheets_queue));
CALL_CL_GUARDED(clReleaseContext, (__sheets_context));
return 0;
}
int main(int argc, char *argv[])
{
int error, xsize, ysize, rgb_max;
int *r, *b, *g;
float *gray, *congray, *congray_cl;
// identity kernel
// float filter[] = {
// 0,0,0,0,0,0,0,
// 0,0,0,0,0,0,0,
// 0,0,0,0,0,0,0,
// 0,0,0,1,0,0,0,
// 0,0,0,0,0,0,0,
// 0,0,0,0,0,0,0,
// 0,0,0,0,0,0,0,
// };
// 45 degree motion blur
float filter[] =
{0, 0, 0, 0, 0, 0.0145, 0,
0, 0, 0, 0, 0.0376, 0.1283, 0.0145,
0, 0, 0, 0.0376, 0.1283, 0.0376, 0,
0, 0, 0.0376, 0.1283, 0.0376, 0, 0,
0, 0.0376, 0.1283, 0.0376, 0, 0, 0,
0.0145, 0.1283, 0.0376, 0, 0, 0, 0,
0, 0.0145, 0, 0, 0, 0, 0};
// mexican hat kernel
// float filter[] = {
// 0, 0,-1,-1,-1, 0, 0,
// 0,-1,-3,-3,-3,-1, 0,
// -1,-3, 0, 7, 0,-3,-1,
// -1,-3, 7,24, 7,-3,-1,
// -1,-3, 0, 7, 0,-3,-1,
// 0,-1,-3,-3,-3,-1, 0,
// 0, 0,-1,-1,-1, 0, 0
// };
if(argc != 3)
{
fprintf(stderr, "Usage: %s image.ppm num_loops\n", argv[0]);
abort();
}
const char* filename = argv[1];
const int num_loops = atoi(argv[2]);
// --------------------------------------------------------------------------
// load image
// --------------------------------------------------------------------------
printf("Reading ``%s''\n", filename);
ppma_read(filename, &xsize, &ysize, &rgb_max, &r, &g, &b);
printf("Done reading ``%s'' of size %dx%d\n", filename, xsize, ysize);
// --------------------------------------------------------------------------
// allocate CPU buffers
// --------------------------------------------------------------------------
posix_memalign((void**)&gray, 32, xsize*ysize*sizeof(float));
if(!gray) { fprintf(stderr, "alloc gray"); abort(); }
posix_memalign((void**)&congray, 32, xsize*ysize*sizeof(float));
if(!congray) { fprintf(stderr, "alloc gray"); abort(); }
posix_memalign((void**)&congray_cl, 32, xsize*ysize*sizeof(float));
if(!congray_cl) { fprintf(stderr, "alloc gray"); abort(); }
// --------------------------------------------------------------------------
// convert image to grayscale
// --------------------------------------------------------------------------
for(int n = 0; n < xsize*ysize; ++n)
gray[n] = (0.21f*r[n])/rgb_max + (0.72f*g[n])/rgb_max + (0.07f*b[n])/rgb_max;
// --------------------------------------------------------------------------
// execute filter on cpu
// --------------------------------------------------------------------------
for(int i = HALF_FILTER_WIDTH; i < ysize - HALF_FILTER_WIDTH; ++i)
{
for(int j = HALF_FILTER_WIDTH; j < xsize - HALF_FILTER_WIDTH; ++j)
{
float sum = 0;
for(int k = -HALF_FILTER_WIDTH; k <= HALF_FILTER_WIDTH; ++k)
{
for(int l = -HALF_FILTER_WIDTH; l <= HALF_FILTER_WIDTH; ++l)
{
sum += gray[(i+k)*xsize + (j+l)] *
filter[(k+HALF_FILTER_WIDTH)*FILTER_WIDTH + (l+HALF_FILTER_WIDTH)];
}
}
congray[i*xsize + j] = sum;
}
}
// --------------------------------------------------------------------------
// output cpu filtered image
// --------------------------------------------------------------------------
printf("Writing cpu filtered image\n");
for(int n = 0; n < xsize*ysize; ++n)
r[n] = g[n] = b[n] = (int)(congray[n] * rgb_max);
error = ppma_write("output_cpu.ppm", xsize, ysize, r, g, b);
if(error) { fprintf(stderr, "error writing image"); abort(); }
// --------------------------------------------------------------------------
// get an OpenCL context and queue
// --------------------------------------------------------------------------
cl_context ctx;
cl_command_queue queue;
create_context_on(CHOOSE_INTERACTIVELY, CHOOSE_INTERACTIVELY, 0, &ctx, &queue, 0);
print_device_info_from_queue(queue);
// --------------------------------------------------------------------------
// load kernels
// --------------------------------------------------------------------------
char *knl_text = read_file("convolution.cl");
cl_kernel knl = kernel_from_string(ctx, knl_text, "convolution", NULL);
free(knl_text);
#ifdef NON_OPTIMIZED
int deviceWidth = xsize;
#else
int deviceWidth = ((xsize + WGX - 1)/WGX)* WGX;
#endif
int deviceHeight = ysize;
size_t deviceDataSize = deviceHeight*deviceWidth*sizeof(float);
// --------------------------------------------------------------------------
// allocate device memory
// --------------------------------------------------------------------------
cl_int status;
cl_mem buf_gray = clCreateBuffer(ctx, CL_MEM_READ_ONLY,
deviceDataSize, 0, &status);
CHECK_CL_ERROR(status, "clCreateBuffer");
cl_mem buf_congray = clCreateBuffer(ctx, CL_MEM_WRITE_ONLY,
deviceDataSize, 0, &status);
CHECK_CL_ERROR(status, "clCreateBuffer");
cl_mem buf_filter = clCreateBuffer(ctx, CL_MEM_READ_ONLY,
FILTER_WIDTH*FILTER_WIDTH*sizeof(float), 0, &status);
CHECK_CL_ERROR(status, "clCreateBuffer");
// --------------------------------------------------------------------------
// transfer to device
// --------------------------------------------------------------------------
#ifdef NON_OPTIMIZED
CALL_CL_SAFE(clEnqueueWriteBuffer(
queue, buf_gray, /*blocking*/ CL_TRUE, /*offset*/ 0,
deviceDataSize, gray, 0, NULL, NULL));
#else
size_t buffer_origin[3] = {0,0,0};
size_t host_origin[3] = {0,0,0};
size_t region[3] = {deviceWidth*sizeof(float), ysize, 1};
clEnqueueWriteBufferRect(queue, buf_gray, CL_TRUE,
buffer_origin, host_origin, region,
deviceWidth*sizeof(float), 0, xsize*sizeof(float), 0,
gray, 0, NULL, NULL);
#endif
CALL_CL_SAFE(clEnqueueWriteBuffer(
queue, buf_filter, /*blocking*/ CL_TRUE, /*offset*/ 0,
FILTER_WIDTH*FILTER_WIDTH*sizeof(float), filter, 0, NULL, NULL));
// --------------------------------------------------------------------------
// run code on device
// --------------------------------------------------------------------------
cl_int rows = ysize;
cl_int cols = xsize;
cl_int filterWidth = FILTER_WIDTH;
cl_int paddingPixels = 2*HALF_FILTER_WIDTH;
size_t local_size[] = { WGX, WGY };
size_t global_size[] = {
((xsize-paddingPixels + local_size[0] - 1)/local_size[0])* local_size[0],
((ysize-paddingPixels + local_size[1] - 1)/local_size[1])* local_size[1],
};
cl_int localWidth = local_size[0] + paddingPixels;
cl_int localHeight = local_size[1] + paddingPixels;
size_t localMemSize = localWidth * localHeight * sizeof(float);
CALL_CL_SAFE(clSetKernelArg(knl, 0, sizeof(buf_gray), &buf_gray));
CALL_CL_SAFE(clSetKernelArg(knl, 1, sizeof(buf_congray), &buf_congray));
CALL_CL_SAFE(clSetKernelArg(knl, 2, sizeof(buf_filter), &buf_filter));
CALL_CL_SAFE(clSetKernelArg(knl, 3, sizeof(rows), &rows));
CALL_CL_SAFE(clSetKernelArg(knl, 4, sizeof(cols), &cols));
CALL_CL_SAFE(clSetKernelArg(knl, 5, sizeof(filterWidth), &filterWidth));
CALL_CL_SAFE(clSetKernelArg(knl, 6, localMemSize, NULL));
CALL_CL_SAFE(clSetKernelArg(knl, 7, sizeof(localHeight), &localHeight));
CALL_CL_SAFE(clSetKernelArg(knl, 8, sizeof(localWidth), &localWidth));
// --------------------------------------------------------------------------
// print kernel info
// --------------------------------------------------------------------------
print_kernel_info(queue, knl);
CALL_CL_SAFE(clFinish(queue));
timestamp_type tic, toc;
get_timestamp(&tic);
for(int loop = 0; loop < num_loops; ++loop)
{
CALL_CL_SAFE(clEnqueueNDRangeKernel(queue, knl, 2, NULL,
global_size, local_size, 0, NULL, NULL));
// Edit: Copy the blurred image to input buffer
#ifdef NON_OPTIMIZED
CALL_CL_SAFE(clEnqueueCopyBuffer(queue, buf_congray, buf_gray, 0, 0,
deviceDataSize, 0, NULL, NULL));
#else
clEnqueueCopyBufferRect(queue, buf_congray, buf_gray,
buffer_origin, host_origin, region,
deviceWidth*sizeof(float), 0,
xsize*sizeof(float), 0,
0, NULL, NULL);
#endif
}
CALL_CL_SAFE(clFinish(queue));
get_timestamp(&toc);
double elapsed = timestamp_diff_in_seconds(tic,toc)/num_loops;
printf("%f s\n", elapsed);
printf("%f MPixels/s\n", xsize*ysize/1e6/elapsed);
printf("%f GBit/s\n", 2*xsize*ysize*sizeof(float)/1e9/elapsed);
printf("%f GFlop/s\n", (xsize-HALF_FILTER_WIDTH)*(ysize-HALF_FILTER_WIDTH)
*FILTER_WIDTH*FILTER_WIDTH/1e9/elapsed);
// --------------------------------------------------------------------------
// transfer back & check
// --------------------------------------------------------------------------
#ifdef NON_OPTIMIZED
CALL_CL_SAFE(clEnqueueReadBuffer(
queue, buf_congray, /*blocking*/ CL_TRUE, /*offset*/ 0,
xsize * ysize * sizeof(float), congray_cl,
0, NULL, NULL));
#else
buffer_origin[0] = 3*sizeof(float);
buffer_origin[1] = 3;
buffer_origin[2] = 0;
host_origin[0] = 3*sizeof(float);
host_origin[1] = 3;
host_origin[2] = 0;
region[0] = (xsize-paddingPixels)*sizeof(float);
region[1] = (ysize-paddingPixels);
region[2] = 1;
clEnqueueReadBufferRect(queue, buf_congray, CL_TRUE,
buffer_origin, host_origin, region,
deviceWidth*sizeof(float), 0, xsize*sizeof(float), 0,
congray_cl, 0, NULL, NULL);
#endif
// --------------------------------------------------------------------------
// output OpenCL filtered image
// --------------------------------------------------------------------------
printf("Writing OpenCL filtered image\n");
// Edit: Keep pixel value in the interval [0, 255] to reduce boundary effect
for(int n = 0; n < xsize*ysize; ++n) {
int color = (int)(congray_cl[n] * rgb_max);
if (color < 0) {
color = 0;
} else if (color > 255) {
color = 255;
}
r[n] = g[n] = b[n] = color;
}
error = ppma_write("output_cl.ppm", xsize, ysize, r, g, b);
if(error) { fprintf(stderr, "error writing image"); abort(); }
// --------------------------------------------------------------------------
// clean up
// --------------------------------------------------------------------------
CALL_CL_SAFE(clReleaseMemObject(buf_congray));
CALL_CL_SAFE(clReleaseMemObject(buf_gray));
CALL_CL_SAFE(clReleaseMemObject(buf_filter));
CALL_CL_SAFE(clReleaseKernel(knl));
CALL_CL_SAFE(clReleaseCommandQueue(queue));
CALL_CL_SAFE(clReleaseContext(ctx));
free(gray);
free(congray);
free(congray_cl);
free(r);
free(b);
free(g);
}
int main (int argc, char *argv[])
{
double *a, *a_reduced;
if (argc != 3)
{
fprintf(stderr, "Usage: %s N nloops\n", argv[0]);
abort();
}
const cl_long N = (cl_long) atol(argv[1]);
const int nloops = atoi(argv[2]);
cl_long Ngroups = (N + LDIM - 1)/LDIM;
Ngroups = (Ngroups + 8 - 1)/8;
cl_context ctx;
cl_command_queue queue;
create_context_on(CHOOSE_INTERACTIVELY, CHOOSE_INTERACTIVELY, 0, &ctx, &queue, 0);
print_device_info_from_queue(queue);
// --------------------------------------------------------------------------
// load kernels
// --------------------------------------------------------------------------
char *knl_text = read_file("full_reduction.cl");
cl_kernel knl = kernel_from_string(ctx, knl_text, "reduction",
"-DLDIM=" STRINGIFY(LDIM));
free(knl_text);
// --------------------------------------------------------------------------
// allocate and initialize CPU memory
// --------------------------------------------------------------------------
posix_memalign((void**)&a, 32, N*sizeof(double));
if (!a) { fprintf(stderr, "alloc a"); abort(); }
posix_memalign((void**)&a_reduced, 32, Ngroups*sizeof(double));
if (!a_reduced) { fprintf(stderr, "alloc a_reduced"); abort(); }
srand48(8);
for(cl_long n = 0; n < N; ++n)
a[n] = (double)drand48();
// a[n] = n;
// --------------------------------------------------------------------------
// allocate device memory
// --------------------------------------------------------------------------
cl_int status;
cl_mem buf_a = clCreateBuffer(ctx, CL_MEM_READ_WRITE, N*sizeof(double),
0, &status);
CHECK_CL_ERROR(status, "clCreateBuffer");
cl_mem buf_a_reduced[2];
buf_a_reduced[0] = clCreateBuffer(ctx, CL_MEM_READ_WRITE,
Ngroups*sizeof(double), 0, &status);
buf_a_reduced[1] = clCreateBuffer(ctx, CL_MEM_READ_WRITE,
Ngroups*sizeof(double), 0, &status);
CHECK_CL_ERROR(status, "clCreateBuffer");
// --------------------------------------------------------------------------
// transfer to device
// --------------------------------------------------------------------------
CALL_CL_SAFE(clEnqueueWriteBuffer(
queue, buf_a, /*blocking*/ CL_TRUE, /*offset*/ 0,
N*sizeof(double), a,
0, NULL, NULL));
timestamp_type tic, toc;
double elapsed;
// --------------------------------------------------------------------------
// run reduction_simple on device
// --------------------------------------------------------------------------
printf("Simple Reduction\n");
double sum_gpu = 0.0;
CALL_CL_SAFE(clFinish(queue));
get_timestamp(&tic);
for(int loop = 0; loop < nloops; ++loop)
{
int r = 0;
size_t Ngroups_loop = Ngroups;
SET_3_KERNEL_ARGS(knl, N, buf_a, buf_a_reduced[r]);
size_t local_size[] = { LDIM };
size_t global_size[] = { Ngroups_loop*LDIM };
CALL_CL_SAFE(clEnqueueNDRangeKernel(queue, knl, 1, NULL,
global_size, local_size, 0, NULL, NULL));
while(Ngroups_loop > 1)
{
cl_long N_reduce = Ngroups_loop;
Ngroups_loop = (N_reduce + LDIM - 1)/LDIM;
Ngroups_loop = (Ngroups_loop + 8 - 1)/8;
size_t local_size[] = { LDIM };
size_t global_size[] = { Ngroups_loop*LDIM };
SET_3_KERNEL_ARGS(knl, N_reduce, buf_a_reduced[r], buf_a_reduced[(r+1)%2]);
CALL_CL_SAFE(clEnqueueNDRangeKernel(queue, knl, 1, NULL,
global_size, local_size, 0, NULL, NULL));
r = (r+1)%2;
}
CALL_CL_SAFE(clEnqueueReadBuffer(
queue, buf_a_reduced[r], /*blocking*/ CL_TRUE, /*offset*/ 0,
Ngroups_loop*sizeof(double), a_reduced, 0, NULL, NULL));
sum_gpu = 0.0;
for(cl_long n = 0; n < Ngroups_loop; ++n)
sum_gpu += a_reduced[n];
}
CALL_CL_SAFE(clFinish(queue));
get_timestamp(&toc);
elapsed = timestamp_diff_in_seconds(tic,toc)/nloops;
printf("%f s\n", elapsed);
printf("%f GB/s\n", N*sizeof(double)/1e9/elapsed);
double sum_cpu = 0.0;
for(cl_long n = 0; n < N; ++n)
sum_cpu += a[n];
printf("Sum CPU: %e\n", sum_cpu);
printf("Sum GPU: %e\n", sum_gpu);
printf("Relative Error: %e\n", fabs(sum_cpu-sum_gpu)/sum_gpu);
// --------------------------------------------------------------------------
// clean up
// --------------------------------------------------------------------------
CALL_CL_SAFE(clReleaseMemObject(buf_a));
CALL_CL_SAFE(clReleaseMemObject(buf_a_reduced[0]));
CALL_CL_SAFE(clReleaseMemObject(buf_a_reduced[1]));
CALL_CL_SAFE(clReleaseKernel(knl));
CALL_CL_SAFE(clReleaseCommandQueue(queue));
CALL_CL_SAFE(clReleaseContext(ctx));
free(a);
free(a_reduced);
return 0;
}
int main (int argc, char *argv[])
{
double *a, *b, *c;
if (argc != 3)
{
fprintf(stderr, "Usage: %s size_of_vector num_adds\n", argv[0]);
abort();
}
const cl_long N = (cl_long) atol(argv[1]);
const int num_adds = atoi(argv[2]);
cl_context ctx;
cl_command_queue queue;
create_context_on(CHOOSE_INTERACTIVELY, CHOOSE_INTERACTIVELY, 0, &ctx, &queue, 0);
print_device_info_from_queue(queue);
// --------------------------------------------------------------------------
// load kernels
// --------------------------------------------------------------------------
char *knl_text = read_file("vec-add-kernel.cl");
cl_kernel knl = kernel_from_string(ctx, knl_text, "sum", NULL);
free(knl_text);
// --------------------------------------------------------------------------
// allocate and initialize CPU memory
// --------------------------------------------------------------------------
posix_memalign((void**)&a, 32, N*sizeof(double));
if (!a) { fprintf(stderr, "alloc a"); abort(); }
posix_memalign((void**)&b, 32, N*sizeof(double));
if (!b) { fprintf(stderr, "alloc b"); abort(); }
posix_memalign((void**)&c, 32, N*sizeof(double));
if (!c) { fprintf(stderr, "alloc c"); abort(); }
for(cl_long n = 0; n < N; ++n)
{
a[n] = n;
b[n] = 2*n;
}
// --------------------------------------------------------------------------
// allocate device memory
// --------------------------------------------------------------------------
cl_int status;
cl_mem buf_a = clCreateBuffer(ctx, CL_MEM_READ_WRITE,
sizeof(double) * N, 0, &status);
CHECK_CL_ERROR(status, "clCreateBuffer");
cl_mem buf_b = clCreateBuffer(ctx, CL_MEM_READ_WRITE,
sizeof(double) * N, 0, &status);
CHECK_CL_ERROR(status, "clCreateBuffer");
cl_mem buf_c = clCreateBuffer(ctx, CL_MEM_READ_WRITE,
sizeof(double) * N, 0, &status);
CHECK_CL_ERROR(status, "clCreateBuffer");
// --------------------------------------------------------------------------
// transfer to device
// --------------------------------------------------------------------------
CALL_CL_SAFE(clEnqueueWriteBuffer(
queue, buf_a, /*blocking*/ CL_TRUE, /*offset*/ 0,
N * sizeof(double), a,
0, NULL, NULL));
CALL_CL_SAFE(clEnqueueWriteBuffer(
queue, buf_b, /*blocking*/ CL_TRUE, /*offset*/ 0,
N * sizeof(double), b,
0, NULL, NULL));
// --------------------------------------------------------------------------
// run code on device
// --------------------------------------------------------------------------
CALL_CL_SAFE(clFinish(queue));
timestamp_type tic, toc;
get_timestamp(&tic);
for(int add = 0; add < num_adds; ++add)
{
SET_4_KERNEL_ARGS(knl, N, buf_a, buf_b, buf_c);
size_t local_size[] = { 128 };
size_t global_size[] = { ((N + local_size[0] - 1)/local_size[0])*
local_size[0] };
CALL_CL_SAFE(clEnqueueNDRangeKernel(queue, knl, 1, NULL,
global_size, local_size, 0, NULL, NULL));
}
CALL_CL_SAFE(clFinish(queue));
get_timestamp(&toc);
double elapsed = timestamp_diff_in_seconds(tic,toc)/num_adds;
printf("%f s\n", elapsed);
printf("%f GB/s\n", 3*N*sizeof(double)/1e9/elapsed);
// --------------------------------------------------------------------------
// transfer back & check
// --------------------------------------------------------------------------
CALL_CL_SAFE(clEnqueueReadBuffer(
queue, buf_c, /*blocking*/ CL_TRUE, /*offset*/ 0,
N * sizeof(double), c,
0, NULL, NULL));
for(cl_long i = 0; i < N; ++i)
if(c[i] != 3*i)
{
printf("BAD %ld\n", (long)i);
abort();
}
printf("GOOD\n");
// --------------------------------------------------------------------------
// clean up
// --------------------------------------------------------------------------
CALL_CL_SAFE(clReleaseMemObject(buf_a));
CALL_CL_SAFE(clReleaseMemObject(buf_b));
CALL_CL_SAFE(clReleaseMemObject(buf_c));
CALL_CL_SAFE(clReleaseKernel(knl));
CALL_CL_SAFE(clReleaseCommandQueue(queue));
CALL_CL_SAFE(clReleaseContext(ctx));
free(a);
free(b);
free(c);
return 0;
}
int main()
{
int enable_profiling = 0;
#ifdef DO_TIMING
enable_profiling = 1;
#endif
//print_platforms_devices();
cl_context ctx;
cl_command_queue queue;
create_context_on("NVIDIA", NULL, 0, &ctx, &queue, enable_profiling);
// --------------------------------------------------------------------------
// load kernels
// --------------------------------------------------------------------------
// read the cl file
char buf[100];
sprintf(buf, "mg-kernel-ver%d.cl", VERSION);
char *knl_text = read_file(buf);
//get work group dimensions and gflop info.
int wg_dims , wg_x, wg_y, wg_z, z_div, fetch_per_pt, flops_per_pt;
if (sscanf(knl_text, "// workgroup: (%d,%d,%d) z_div:%d fetch_per_pt:%d flops_per_pt:%d",
&wg_x, &wg_y, &wg_z, &z_div, &fetch_per_pt, &flops_per_pt) == 6)
{
wg_dims = 3;
}
else if (sscanf(knl_text, "// workgroup: (%d,%d) fetch_per_pt:%d flops_per_pt:%d",
&wg_x, &wg_y, &fetch_per_pt, &flops_per_pt) == 4)
{
wg_dims = 2;
wg_z = -1;
z_div = -1;
}
else
{
perror("reading workgroup spec");
abort();
}
#ifdef USE_DOUBLE
char *compile_opt = "-DFTYPE=double";
#else
char *compile_opt = "-DFTYPE=float";
#endif
// creation of the kernel
cl_kernel poisson_knl = kernel_from_string(ctx, knl_text, "fd_update", compile_opt);
free(knl_text); // my compiler complains about this one. OJO!!
// --------------------------------------------------------------------------
// set up grid
// --------------------------------------------------------------------------
const unsigned points = POINTS;
const ftype minus_bdry = -1, plus_bdry = 1;
// We're dividing into (points-1) intervals.
ftype dx = (plus_bdry-minus_bdry)/(points-1);
// --------------------------------------------------------------------------
// allocate and initialize CPU memory
// --------------------------------------------------------------------------
int use_alignment;
unsigned dim_other = points; //if order 2 then 1 point extra on each side
#ifdef USE_ALIGNMENT
// adjusts dimension so that the next row starts in a number divisible by 16
unsigned dim_x = ((dim_other + 15) / 16) * 16;
unsigned field_start = 0;
use_alignment = 1;
#else
unsigned dim_x = dim_other;
unsigned field_start = 0;// this one puts me right at the beginning
use_alignment = 0;
#endif
// --------Allocate forcing uexact, r and u vectors -------------------------
const size_t field_size = 0+dim_x*dim_x*dim_x; // extra large to fit the 2^n constrain in GPU
ftype *f = malloc(field_size*sizeof(ftype));
CHECK_SYS_ERROR(!f, "allocating f");
ftype *u = malloc (field_size*sizeof(ftype));
CHECK_SYS_ERROR(!u, "allocating u");
ftype *uexact = malloc (field_size*sizeof(ftype));
CHECK_SYS_ERROR(!uexact, "allocating uexact");
ftype *r = malloc(field_size * sizeof(ftype));
CHECK_SYS_ERROR(!r, "allocating residual r");
// --------------------------------------------------------------------------
// initialize
// --------------------------------------------------------------------------
// zero out (necessary to initialize everything bec. I measure norms)
for (size_t i = 0; i < field_size; ++i){
f[i] = 0;
u[i] = 0;
uexact[i] = 0;
r[i] = 0;
}
// set up the forcing field
init_f (points, f, dx, field_start, dim_x, dim_other, minus_bdry);
// Initialize u with initial boundary conditions
init_u ( points, u , minus_bdry, plus_bdry, dx, field_start, dim_x, dim_other);
// Initialize the exact solution
init_uexact(points, u, uexact, dx, field_size, field_start, dim_x, dim_other);
// --------------------------------------------------------------------------
// Setup the v-cycles
// --------------------------------------------------------------------------
unsigned n1, n2, n3, ncycles;
n1 = 50;
n2 = 60;
n3 = 1;
ncycles = 2;
ftype *sweeps = malloc (ncycles*sizeof(ftype));
ftype *rnorm = malloc (ncycles*sizeof(ftype));
ftype *enorm = malloc (ncycles*sizeof(ftype));
ftype rtol = 1.0e-05;
// Find the norm of the residual (choose your method)
sweeps[0] =0;
resid (r, f, u, dx, field_size, field_start, dim_x, dim_other);
rnorm[0] = norm( r , field_size) * dx;
U_error(u, uexact, r, field_size);
enorm[0] = norm( r, field_size ) * dx;
for(unsigned icycle = 1; icycle <= ncycles; icycle++){
mgv(f, u, dx, n1, n2, n3, field_size, points, use_alignment, dim_x, ctx, queue, poisson_knl, wg_dims , wg_x, wg_y, wg_z, z_div, fetch_per_pt, flops_per_pt); //update u through a v-cycle
sweeps[icycle] = sweeps[icycle -1] + (4 * (n1 + n2)/3);
resid (r, f, u, dx, field_size, field_start, dim_x, dim_other);
rnorm[icycle] = norm( r, field_size ) * dx;
U_error(u, uexact, r, field_size);
enorm[icycle] = norm( r, field_size ) * dx;
//cfacts = (rnorm(icycle)/rnorm(icycle - 1))^(1 / (n1 + n2)) not necessary
//disp something here if I want to.
//printf("norm of the cycle %f", enorm[icycle]);
if(rnorm[icycle] <= rtol * rnorm[0])
break;
}
#ifdef DO_TIMING
printf(" ftype:%d ver:%d align:%d pts:%d\tgflops:%.1f\tmcells:%.1f\tgbytes:%.1f [/sec]\tout_gflops:%.6f\n", (int) sizeof(ftype), VERSION, use_alignment, points, gflops_performed/seconds_taken, mcells_updated/seconds_taken, gbytes_accessed/seconds_taken, gflops_performed/tot_secs);
#endif
// --------------------------------------------------------------------------
// clean up
// --------------------------------------------------------------------------
CALL_CL_GUARDED(clReleaseKernel, (poisson_knl));
CALL_CL_GUARDED(clReleaseCommandQueue, (queue));
CALL_CL_GUARDED(clReleaseContext, (ctx));
}
void main(int argc, char** argv)
{
//int k = atoi(argv[1]);
//int N = pow(2,k);
int N=1024;
int k=10;
float * a = (float *) malloc(sizeof(float)*N* N * 2);
float * b = (float *) malloc(sizeof(float) *N*N * 2);
float * c = (float *) malloc(sizeof(float) * N*N* 2);
float p = 2*M_PI ;
for (int i =0; i< N*N; i++)
{
a[2*i] = 1;
a[2*i+1] = 0;
b[2*i] = 1;
b[2*i+1] = 0;
}
#if 0
srand(1);
for(int i =0;i<N*N;i++)
{
a[2*i]=sin(i%N *2 *M_PI);
//printf("%f\n",uu[2*i]);
a[2*i+1] =0 ;
}
#endif
print_platforms_devices();
cl_context ctx;
cl_command_queue queue;
create_context_on("NVIDIA","GeForce GTX 590",0,&ctx,&queue,0);
cl_int status;
cl_mem buf_a = clCreateBuffer(ctx, CL_MEM_READ_WRITE,
sizeof(float) *N *N* 2 , 0, &status);
CHECK_CL_ERROR(status, "clCreateBuffer");
cl_mem buf_b = clCreateBuffer(ctx, CL_MEM_READ_WRITE,
sizeof(float) * N *N* 2 , 0, &status);
CHECK_CL_ERROR(status, "clCreateBuffer");
cl_mem buf_c = clCreateBuffer(ctx, CL_MEM_READ_WRITE,
sizeof(float) * N *N* 2 , 0, &status);
CHECK_CL_ERROR(status, "clCreateBuffer");
cl_mem buf_d = clCreateBuffer(ctx, CL_MEM_READ_WRITE,
sizeof(float)*N *N* 2 , 0, &status);
CHECK_CL_ERROR(status, "clCreateBuffer");
cl_mem buf_e = clCreateBuffer(ctx, CL_MEM_READ_WRITE,
sizeof(float) *N *N* 2 , 0, &status);
CHECK_CL_ERROR(status, "clCreateBuffer");
cl_mem buf_f = clCreateBuffer(ctx, CL_MEM_READ_WRITE,
sizeof(float) *N *N* 2 , 0, &status);
CHECK_CL_ERROR(status, "clCreateBuffer");
cl_mem buf_g = clCreateBuffer(ctx, CL_MEM_READ_WRITE,
sizeof(float) *N *N* 2 , 0, &status);
CHECK_CL_ERROR(status, "clCreateBuffer");
CALL_CL_GUARDED(clEnqueueWriteBuffer, (
queue, buf_a, /*blocking*/ CL_TRUE, /*offset*/ 0,
sizeof(float) *N*N*2, a,
0, NULL, NULL));
CALL_CL_GUARDED(clEnqueueWriteBuffer, (
queue, buf_b, /*blocking*/ CL_TRUE, /*offset*/ 0,
sizeof(float) *N *N* 2, b,
0, NULL, NULL));
CALL_CL_GUARDED(clEnqueueWriteBuffer, (
queue, buf_c, /*blocking*/ CL_TRUE, /*offset*/ 0,
sizeof(float) *N* N*2, c,
0, NULL, NULL));
char *knl_text = read_file("vec_add.cl");
cl_kernel vec_add = kernel_from_string(ctx, knl_text, "sum", NULL);
free(knl_text);
knl_text = read_file("mat_etr_mul.cl");
cl_kernel mat_etr_mul = kernel_from_string(ctx, knl_text, "mult", NULL);
free(knl_text);
knl_text = read_file("radix-4-float.cl");
cl_kernel fft1D = kernel_from_string(ctx, knl_text, "fft1D", NULL);
free(knl_text);
knl_text = read_file("radix-4-init.cl");
cl_kernel fft_init = kernel_from_string(ctx, knl_text, "fft1D_init", NULL);
free(knl_text);
knl_text = read_file("radix-4-interm.cl");
cl_kernel fft_interm = kernel_from_string(ctx, knl_text, "fft1D", NULL);
free(knl_text);
knl_text = read_file("transpose-soln-gpu.cl");
cl_kernel mat_trans = kernel_from_string(ctx, knl_text, "transpose", NULL);
free(knl_text);
knl_text = read_file("radix-4-modi.cl");
cl_kernel fft_init_w = kernel_from_string(ctx, knl_text, "fft1D_init", NULL);
free(knl_text);
knl_text = read_file("vec_zero.cl");
cl_kernel vec_zero = kernel_from_string(ctx, knl_text, "zero", NULL);
free(knl_text);
knl_text = read_file("reduction.cl");
cl_kernel reduct_mul = kernel_from_string(ctx, knl_text, "reduction_mult", NULL);
free(knl_text);
knl_text = read_file("reduction1D.cl");
cl_kernel reduct = kernel_from_string(ctx, knl_text, "reduction", NULL);
free(knl_text);
knl_text = read_file("reduction-init.cl");
cl_kernel reduct_init = kernel_from_string(ctx, knl_text, "reduction_init", NULL);
free(knl_text);
knl_text = read_file("reduct-energy.cl");
cl_kernel reduct_eng = kernel_from_string(ctx, knl_text, "reduction_eng", NULL);
free(knl_text);
knl_text = read_file("resid.cl");
cl_kernel resid = kernel_from_string(ctx, knl_text, "resid", NULL);
free(knl_text);
knl_text = read_file("resid-init.cl");
cl_kernel resid_init = kernel_from_string(ctx, knl_text, "resid_init", NULL);
free(knl_text);
knl_text = read_file("radix-4-big.cl");
cl_kernel fft_big = kernel_from_string(ctx, knl_text, "fft1D_big", NULL);
free(knl_text);
knl_text = read_file("radix-4-big-clean.cl");
cl_kernel fft_clean = kernel_from_string(ctx, knl_text, "fft1D_clean", NULL);
free(knl_text);
knl_text = read_file("radix-4-2D.cl");
cl_kernel fft_2D = kernel_from_string(ctx, knl_text, "fft2D_big", NULL);
free(knl_text);
knl_text = read_file("radix-4-2D-clean.cl");
cl_kernel fft_2D_clean = kernel_from_string(ctx, knl_text, "fft2D_clean", NULL);
free(knl_text);
knl_text = read_file("mat-trans-3D.cl");
cl_kernel mat_trans_3D = kernel_from_string(ctx, knl_text, "transpose_3D", NULL);
free(knl_text);
int Ns =1 ;
int direction = 1;
timestamp_type time1, time2;
struct parameter param;
param.N = N;
param.epsilon = 0.1;
param.s =1;
float kk =1e-4;
param.h = 2*PI/N;
param.N = N;
param.maxCG = 1000;
param.maxN = 5;
//Minimum and starting time step
float mink = 1e-7;
float startk = 1e-4;
// Tolerances
param.Ntol = 1e-4;
param.cgtol = 1e-7;
float ksafety = 0.8;
float kfact = 1.3;
float kfact2 = 1/1.3;
float Nfact = 0.7;
float CGfact = 0.7;
double elapsed ;
CALL_CL_GUARDED(clFinish, (queue));
get_timestamp(&time1);
//for(int s=0;s<100;s++)
//fft_1D_big(buf_a,buf_b,buf_c,N,fft_big,fft_clean,mat_trans,queue,direction,0);
//fft_1D_new(buf_a,buf_b,buf_c,N,fft_init,fft_interm, fft1D,queue,direction,0);
//fft_1D(buf_a,buf_b,buf_c,N,fft_init, fft1D,queue,direction,0);
//fft2D(buf_a,buf_b,buf_c,buf_d,N,fft_init,fft1D,mat_trans,queue, 1);
//fft2D_new(buf_a,buf_b,buf_c,buf_d,N,fft_init,fft_interm,fft1D,mat_trans,queue, 1);
//fft2D_big(buf_a,buf_b,buf_c,buf_d,N,fft_big,fft_clean,mat_trans,queue,direction);
//fft2D_big_new(buf_a,buf_b,buf_c,buf_d,N,fft_2D,fft_2D_clean,
//mat_trans,mat_trans_3D,queue,direction);
//fft_w(buf_a,buf_b,buf_c,buf_d,buf_e,N,0.1,0,1,fft_init_w,fft_init,fft1D,mat_trans,queue);
#if 0
frhs(buf_a,buf_b,buf_c,buf_d,buf_e,¶m,fft1D_init,fft1D,mat_trans,
vec_add, queue);
#endif
#if 0
float E1 = energy(buf_a, buf_b, buf_c,buf_d, buf_e,buf_f,1e-4,
¶m, fft_init,fft1D,mat_trans,reduct_eng,
reduct,queue);
#endif
//float reside = residual(buf_a,buf_b,resid,resid_init,queue,N*N);
/*fft_d_q(buf_a,buf_b,buf_c,buf_d, N,0.1,k ,1,
fft1D_init,
fft1D,mat_trans,queue);*/
//for(int j= 0;j<N;j++)
//{
//fft_1D_w_orig(buf_a,buf_b,buf_c,N,fft1D_init,fft1D,queue,1,j);
//}
//fft_shar(buf_a,buf_b,buf_c,buf_d,N,0.1,0,1,fft1D_init,fft1D,mat_trans,queue);
//mat__trans(buf_a,buf_b,N,mat_trans,queue,4,0.1,0,1);
//double elapsed = reduction_mult(buf_a, buf_b,buf_c,N*N,reduct_mul,reduct,queue);
CALL_CL_GUARDED(clFinish, (queue));
get_timestamp(&time1);
fft_1D_big(buf_a,buf_b,buf_c,N*N,fft_big,fft_clean,mat_trans,queue,direction,0);
CALL_CL_GUARDED(clFinish, (queue));
get_timestamp(&time2);
elapsed = timestamp_diff_in_seconds(time1,time2);
printf("Hierarchy 1D FFT of size %d array on gpu takes %f s\n", N*N,elapsed);
printf("achieve %f GFLOPS \n",6*2*N*N*k/elapsed*1e-9);
printf("---------------------------------------------\n");
CALL_CL_GUARDED(clFinish, (queue));
get_timestamp(&time1);
fft2D(buf_a,buf_b,buf_c,buf_d,N,fft_init,fft1D,mat_trans,queue, 1);
CALL_CL_GUARDED(clFinish, (queue));
get_timestamp(&time2);
elapsed = timestamp_diff_in_seconds(time1,time2);
printf("Navie 2D FFT of size %d * %d matrix on gpu takes %f s\n", N,N,elapsed);
printf("achieve %f GFLOPS \n",6*2*N*N*k/elapsed*1e-9);
printf("---------------------------------------------\n");
//printf("data access from global achieve %f GB/s\n",sizeof(float)*2*16*N*N/elapsed*1e-9);
CALL_CL_GUARDED(clFinish, (queue));
get_timestamp(&time1);
fft2D_new(buf_a,buf_b,buf_c,buf_d,N,fft_init,fft_interm,fft1D,mat_trans,queue, 1);
CALL_CL_GUARDED(clFinish, (queue));
get_timestamp(&time2);
elapsed = timestamp_diff_in_seconds(time1,time2);
printf("local data exchange 2D FFT of size %d * %d matrix on gpu takes %f s\n", N,N,elapsed);
printf("achieve %f GFLOPS \n",6*2*N*N*k/elapsed*1e-9);
printf("---------------------------------------------\n");
CALL_CL_GUARDED(clFinish, (queue));
get_timestamp(&time1);
fft2D_big(buf_a,buf_b,buf_c,buf_d,N,fft_big,fft_clean,mat_trans,queue,direction);
CALL_CL_GUARDED(clFinish, (queue));
get_timestamp(&time2);
elapsed = timestamp_diff_in_seconds(time1,time2);
printf("Hierarchy 2D FFT of size %d * %d matrix on gpu takes %f s\n", N,N,elapsed);
printf("achieve %f GFLOPS \n",6*2*N*N*k/elapsed*1e-9);
printf("---------------------------------------------\n");
CALL_CL_GUARDED(clFinish, (queue));
get_timestamp(&time1);
fft2D_big_new(buf_a,buf_b,buf_c,buf_d,N,fft_2D,fft_2D_clean,
mat_trans,mat_trans_3D,queue,direction);
CALL_CL_GUARDED(clFinish, (queue));
get_timestamp(&time2);
elapsed = timestamp_diff_in_seconds(time1,time2);
printf("Using 2D kernel 2D FFT of size %d * %d matrix on gpu takes %f s\n", N,N,elapsed);
printf("achieve %f GFLOPS \n",6*2*N*N*k/elapsed*1e-9);
printf("---------------------------------------------\n");
get_timestamp(&time1);
direction = -1;
//fft_1D(buf_b,buf_c,buf_d,N,fft_init, fft1D,queue,direction,0);
fft2D(buf_b,buf_c,buf_d,buf_e,N,fft_init,fft1D,mat_trans,queue, direction);
//fft2D_new(buf_b,buf_c,buf_e,buf_d,N,fft_init,fft_interm,fft1D,mat_trans,queue, -1);
//fft2D_big(buf_b,buf_c,buf_d,buf_e,N,fft_big,fft_clean,mat_trans,queue,direction);
CALL_CL_GUARDED(clFinish, (queue));
get_timestamp(&time2);
elapsed = timestamp_diff_in_seconds(time1,time2);
//printf("1D inverse %f s\n", elapsed);
#if 0
float test;
CALL_CL_GUARDED(clFinish, (queue));
CALL_CL_GUARDED(clEnqueueReadBuffer, (
queue, buf_b, /*blocking*/ CL_TRUE, /*offset*/ 0,
sizeof(float), &test,
0, NULL, NULL));
printf("test success and %f \n",test);
#endif
#if 0
CALL_CL_GUARDED(clFinish, (queue));
CALL_CL_GUARDED(clEnqueueReadBuffer, (
queue, buf_c, /*blocking*/ CL_TRUE, /*offset*/ 0,
2*N*N* sizeof(float), c,
0, NULL, NULL));
/*for(int i =0; i< N; i++)
{
printf("a%f+ i*",a[2*i]);
printf("%f\n",a[2*i+1]);
}*/
int T = 10<N? 10:N ;
for(int i =0; i< T; i++)
{
printf("%f + i*",a[2*i]);
printf("%f\t",a[2*i+1]);
printf("%f + i*",c[2*i]);
printf("%f\n",c[2*i+1]);
}
#endif
/* for( Ns = 1;Ns < N; Ns *= 2 )
{
for (int j = 0; j<N/2; j++)
{
fftiteration(j,N,Ns,a,b);
}
float * d;
d = a ;
a = b;
b = d;
//printf("ok\n");
}
*/
CALL_CL_GUARDED(clReleaseMemObject, (buf_a));
CALL_CL_GUARDED(clReleaseMemObject, (buf_b));
CALL_CL_GUARDED(clReleaseMemObject, (buf_c));
CALL_CL_GUARDED(clReleaseMemObject, (buf_d));
CALL_CL_GUARDED(clReleaseMemObject, (buf_e));
CALL_CL_GUARDED(clReleaseKernel, (fft1D));
CALL_CL_GUARDED(clReleaseKernel, (fft_init));
CALL_CL_GUARDED(clReleaseKernel, (vec_add));
CALL_CL_GUARDED(clReleaseKernel, (reduct_mul));
CALL_CL_GUARDED(clReleaseKernel, (reduct));
CALL_CL_GUARDED(clReleaseKernel, (mat_trans));
CALL_CL_GUARDED(clReleaseCommandQueue, (queue));
CALL_CL_GUARDED(clReleaseContext, (ctx));
}
