int main(int argc, char *argv[]){ int size = 18; long long us = 0; stopwatch sw; int opt, option_index=0; int status = 1; if(argc < 2){ fprintf(stderr, "Usage: %s [-s board_size]\n", argv[0]); exit(EXIT_FAILURE); } while ((opt = getopt_long(argc, argv, "s:", long_options, &option_index)) != -1 ) { if(opt == 's')size = atoi(optarg); else{ fprintf(stderr, "Usage: %s [-s board_size]\n", argv[0]); exit(EXIT_FAILURE); } } stopwatch_start(&sw); long long solutions = nqueen_cpu(size, &us); stopwatch_stop(&sw); fprintf(stderr, "The number of solutions is %lld, the number of unique solutions is " "%lld and the total time it took is %lf seconds\n", solutions, us, get_interval_by_sec(&sw)); printf("{ \"status\": %d, \"options\": \"-s %d\", \"time\": %f, \"output\": \"[%lld, %lld]\" }\n", status, size, get_interval_by_sec(&sw), solutions, us); return 0; }
int main(int argc, char* argv[]) { stopwatch sw; if (argc < 2) { printf("usage: %s x\n", argv[0]); return 0; } int n = atoi(argv[1]); stopwatch_start(&sw); int x = fib(n); stopwatch_stop(&sw); printf("{ \"time\": %f, \"output\": %d }\n", get_interval_by_sec(&sw), x); }
int main(int argc, char **argv) { struct stopwatch sw; if (argc < 2) { printf("usage: %s <input size>\n", argv[0]); return 1; } int n = atoi(argv[1]); stopwatch_start(&sw); int x = numprime(n); stopwatch_stop(&sw); printf("{ \"time\": %f, \"result\": %d }\n", get_interval_by_sec(&sw), x); return 0; }
int main(int argc, char** argv) { cl_int err; int usegpu = USEGPU; int do_verify = 0; int opt, option_index=0; unsigned int correct; size_t global_size; size_t local_size; cl_device_id device_id; cl_context context; cl_command_queue commands; cl_program program; cl_kernel kernel; stopwatch sw; cl_mem csr_ap; cl_mem csr_aj; cl_mem csr_ax; cl_mem x_loc; cl_mem y_loc; FILE *kernelFile; char *kernelSource; size_t kernelLength; size_t lengthRead; ocd_init(&argc, &argv, NULL); ocd_options opts = ocd_get_options(); platform_id = opts.platform_id; n_device = opts.device_id; while ((opt = getopt_long(argc, argv, "::vc::", long_options, &option_index)) != -1 ) { switch(opt){ //case 'i': //input_file = optarg; //break; case 'v': fprintf(stderr, "verify\n"); do_verify = 1; break; case 'c': fprintf(stderr, "using cpu\n"); usegpu = 0; break; default: fprintf(stderr, "Usage: %s [-v Warning: lots of output] [-c use CPU]\n", argv[0]); exit(EXIT_FAILURE); } } /* Fill input set with random float values */ int i; csr_matrix csr; csr = laplacian_5pt(512); int k = 0; for(k = 0; k < csr.num_nonzeros; k++){ csr.Ax[k] = 1.0 - 2.0 * (rand() / (RAND_MAX + 1.0)); } //The other arrays float * x_host = float_new_array(csr.num_cols); float * y_host = float_new_array(csr.num_rows); unsigned int ii; for(ii = 0; ii < csr.num_cols; ii++){ x_host[ii] = rand() / (RAND_MAX + 1.0); } for(ii = 0; ii < csr.num_rows; ii++){ y_host[ii] = rand() / (RAND_MAX + 2.0); } /* Retrieve an OpenCL platform */ device_id = GetDevice(platform_id, n_device); /* Create a compute context */ context = clCreateContext(0, 1, &device_id, NULL, NULL, &err); CHKERR(err, "Failed to create a compute context!"); /* Create a command queue */ commands = clCreateCommandQueue(context, device_id, CL_QUEUE_PROFILING_ENABLE, &err); CHKERR(err, "Failed to create a command queue!"); /* Load kernel source */ kernelFile = fopen("spmv_csr_kernel.cl", "r"); fseek(kernelFile, 0, SEEK_END); kernelLength = (size_t) ftell(kernelFile); kernelSource = (char *) malloc(sizeof(char)*kernelLength); rewind(kernelFile); lengthRead = fread((void *) kernelSource, kernelLength, 1, kernelFile); fclose(kernelFile); /* Create the compute program from the source buffer */ program = clCreateProgramWithSource(context, 1, (const char **) &kernelSource, &kernelLength, &err); CHKERR(err, "Failed to create a compute program!"); /* Free kernel source */ free(kernelSource); /* Build the program executable */ err = clBuildProgram(program, 0, NULL, NULL, NULL, NULL); if (err == CL_BUILD_PROGRAM_FAILURE) { char *buildLog; size_t logLen; err = clGetProgramBuildInfo(program, device_id, CL_PROGRAM_BUILD_LOG, 0, NULL, &logLen); buildLog = (char *) malloc(sizeof(char)*logLen); err = clGetProgramBuildInfo(program, device_id, CL_PROGRAM_BUILD_LOG, logLen, (void *) buildLog, NULL); fprintf(stderr, "CL Error %d: Failed to build program! Log:\n%s", err, buildLog); free(buildLog); exit(1); } CHKERR(err, "Failed to build program!"); /* Create the compute kernel in the program we wish to run */ kernel = clCreateKernel(program, "csr", &err); CHKERR(err, "Failed to create a compute kernel!"); /* Create the input and output arrays in device memory for our calculation */ csr_ap = clCreateBuffer(context, CL_MEM_READ_ONLY, sizeof(unsigned int)*csr.num_rows+4, NULL, &err); CHKERR(err, "Failed to allocate device memory!"); csr_aj = clCreateBuffer(context, CL_MEM_READ_ONLY, sizeof(unsigned int)*csr.num_nonzeros, NULL, &err); CHKERR(err, "Failed to allocate device memory!"); csr_ax = clCreateBuffer(context, CL_MEM_READ_ONLY, sizeof(float)*csr.num_nonzeros, NULL, &err); CHKERR(err, "Failed to allocate device memory!"); x_loc = clCreateBuffer(context, CL_MEM_READ_ONLY, sizeof(float)*csr.num_cols, NULL, &err); CHKERR(err, "Failed to allocate device memory!"); y_loc = clCreateBuffer(context, CL_MEM_READ_ONLY, sizeof(float)*csr.num_rows, NULL, &err); CHKERR(err, "Failed to allocate device memory!"); /* beginning of timing point */ stopwatch_start(&sw); /* Write our data set into the input array in device memory */ err = clEnqueueWriteBuffer(commands, csr_ap, CL_TRUE, 0, sizeof(unsigned int)*csr.num_rows+4, csr.Ap, 0, NULL, &ocdTempEvent); clFinish(commands); START_TIMER(ocdTempEvent, OCD_TIMER_H2D, "CSR Data Copy", ocdTempTimer) END_TIMER(ocdTempTimer) CHKERR(err, "Failed to write to source array!"); err = clEnqueueWriteBuffer(commands, csr_aj, CL_TRUE, 0, sizeof(unsigned int)*csr.num_nonzeros, csr.Aj, 0, NULL, &ocdTempEvent); clFinish(commands); START_TIMER(ocdTempEvent, OCD_TIMER_H2D, "CSR Data Copy", ocdTempTimer) END_TIMER(ocdTempTimer) CHKERR(err, "Failed to write to source array!"); err = clEnqueueWriteBuffer(commands, csr_ax, CL_TRUE, 0, sizeof(float)*csr.num_nonzeros, csr.Ax, 0, NULL, &ocdTempEvent); clFinish(commands); START_TIMER(ocdTempEvent, OCD_TIMER_H2D, "CSR Data Copy", ocdTempTimer) END_TIMER(ocdTempTimer) CHKERR(err, "Failed to write to source array!"); err = clEnqueueWriteBuffer(commands, x_loc, CL_TRUE, 0, sizeof(float)*csr.num_cols, x_host, 0, NULL, &ocdTempEvent); clFinish(commands); START_TIMER(ocdTempEvent, OCD_TIMER_H2D, "CSR Data Copy", ocdTempTimer) END_TIMER(ocdTempTimer) CHKERR(err, "Failed to write to source array!"); err = clEnqueueWriteBuffer(commands, y_loc, CL_TRUE, 0, sizeof(float)*csr.num_rows, y_host, 0, NULL, &ocdTempEvent); clFinish(commands); START_TIMER(ocdTempEvent, OCD_TIMER_H2D, "CSR Data Copy", ocdTempTimer) CHKERR(err, "Failed to write to source array!"); END_TIMER(ocdTempTimer) /* Set the arguments to our compute kernel */ err = 0; err = clSetKernelArg(kernel, 0, sizeof(unsigned int), &csr.num_rows); err |= clSetKernelArg(kernel, 1, sizeof(cl_mem), &csr_ap); err |= clSetKernelArg(kernel, 2, sizeof(cl_mem), &csr_aj); err |= clSetKernelArg(kernel, 3, sizeof(cl_mem), &csr_ax); err |= clSetKernelArg(kernel, 4, sizeof(cl_mem), &x_loc); err |= clSetKernelArg(kernel, 5, sizeof(cl_mem), &y_loc); CHKERR(err, "Failed to set kernel arguments!"); /* Get the maximum work group size for executing the kernel on the device */ err = clGetKernelWorkGroupInfo(kernel, device_id, CL_KERNEL_WORK_GROUP_SIZE, sizeof(size_t), (void *) &local_size, NULL); CHKERR(err, "Failed to retrieve kernel work group info!"); /* Execute the kernel over the entire range of our 1d input data set */ /* using the maximum number of work group items for this device */ global_size = csr.num_rows; err = clEnqueueNDRangeKernel(commands, kernel, 1, NULL, &global_size, &local_size, 0, NULL, &ocdTempEvent); clFinish(commands); START_TIMER(ocdTempEvent, OCD_TIMER_KERNEL, "CSR Kernel", ocdTempTimer) END_TIMER(ocdTempTimer) CHKERR(err, "Failed to execute kernel!"); /* Wait for the command commands to get serviced before reading back results */ float output[csr.num_rows]; /* Read back the results from the device to verify the output */ err = clEnqueueReadBuffer(commands, y_loc, CL_TRUE, 0, sizeof(float)*csr.num_rows, output, 0, NULL, &ocdTempEvent); clFinish(commands); START_TIMER(ocdTempEvent, OCD_TIMER_D2H, "CSR Data Copy", ocdTempTimer) END_TIMER(ocdTempTimer) CHKERR(err, "Failed to read output array!"); /* end of timing point */ stopwatch_stop(&sw); printf("Time consumed(ms): %lf Gflops: %f \n", 1000*get_interval_by_sec(&sw), (2.0 * (double) csr.num_nonzeros / get_interval_by_sec(&sw)) / 1e9); /* Validate our results */ if(do_verify){ for (i = 0; i < csr.num_rows; i++){ printf("row: %d output: %f \n", i, output[i]); } } int row = 0; float sum = 0; int row_start = 0; int row_end = 0; for(row =0; row < csr.num_rows; row++){ sum = y_host[row]; row_start = csr.Ap[row]; row_end = csr.Ap[row+1]; unsigned int jj = 0; for (jj = row_start; jj < row_end; jj++){ sum += csr.Ax[jj] * x_host[csr.Aj[jj]]; } y_host[row] = sum; } for (i = 0; i < csr.num_rows; i++){ if((fabsf(y_host[i]) - fabsf(output[i])) > .001) printf("Possible error, difference greater then .001 at row %d \n", i); } /* Print a brief summary detailing the results */ ocd_finalize(); /* Shutdown and cleanup */ clReleaseMemObject(csr_ap); clReleaseMemObject(csr_aj); clReleaseMemObject(csr_ax); clReleaseMemObject(x_loc); clReleaseMemObject(y_loc); clReleaseProgram(program); clReleaseKernel(kernel); clReleaseCommandQueue(commands); clReleaseContext(context); return 0; }
int main ( int argc, char *argv[] ) { int matrix_dim = 32; /* default matrix_dim */ int opt, option_index=0; func_ret_t ret; const char *input_file = NULL; float *m, *mm; stopwatch sw; cl_device_id clDevice; cl_context clContext; cl_command_queue clCommands; cl_program clProgram; cl_kernel clKernel_diagonal; cl_kernel clKernel_perimeter; cl_kernel clKernel_internal; cl_int dev_type; cl_int errcode; FILE *kernelFile; char *kernelSource; size_t kernelLength; cl_mem d_m; ocd_init(&argc, &argv, NULL); ocd_options opts = ocd_get_options(); platform_id = opts.platform_id; device_id = opts.device_id; while ((opt = getopt_long(argc, argv, "::vs:i:", long_options, &option_index)) != -1 ) { switch(opt) { case 'i': input_file = optarg; break; case 'v': do_verify = 1; break; case 's': matrix_dim = atoi(optarg); fprintf(stderr, "Currently not supported, use -i instead\n"); fprintf(stderr, "Usage: %s [-v] [-s matrix_size|-i input_file|-p platform|-d device]\n", argv[0]); exit(EXIT_FAILURE); case '?': fprintf(stderr, "invalid option\n"); break; case ':': fprintf(stderr, "missing argument\n"); break; default: fprintf(stderr, "Usage: %s [-v] [-s matrix_size|-i input_file||-p platform|-d device]\n", argv[0]); exit(EXIT_FAILURE); } } if ( (optind < argc) || (optind == 1)) { fprintf(stderr, "Usage: %s [-v] [-s matrix_size|-i input_file|-p platform|-d device]\n", argv[0]); exit(EXIT_FAILURE); } if (input_file) { printf("Reading matrix from file %s\n", input_file); ret = create_matrix_from_file(&m, input_file, &matrix_dim); if (ret != RET_SUCCESS) { m = NULL; fprintf(stderr, "error create matrix from file %s\n", input_file); exit(EXIT_FAILURE); } } else { printf("No input file specified!\n"); exit(EXIT_FAILURE); } if (do_verify) { printf("Before LUD\n"); print_matrix(m, matrix_dim); matrix_duplicate(m, &mm, matrix_dim); } // errcode = clGetPlatformIDs(NUM_PLATFORM, clPlatform, NULL); // CHECKERR(errcode); // // errcode = clGetDeviceIDs(clPlatform[PLATFORM_ID], USEGPU ? CL_DEVICE_TYPE_GPU : CL_DEVICE_TYPE_CPU, 1, &clDevice, NULL); // CHECKERR(errcode); #ifdef USEGPU dev_type = CL_DEVICE_TYPE_GPU; #elif defined(USE_AFPGA) dev_type = CL_DEVICE_TYPE_ACCELERATOR; #else dev_type = CL_DEVICE_TYPE_CPU; #endif clDevice = GetDevice(platform_id, device_id,dev_type); size_t max_worksize[3]; errcode = clGetDeviceInfo(clDevice, CL_DEVICE_MAX_WORK_ITEM_SIZES,sizeof(size_t)*3, &max_worksize, NULL); CHECKERR(errcode); while(BLOCK_SIZE*BLOCK_SIZE>max_worksize[0]) BLOCK_SIZE = BLOCK_SIZE/2; clContext = clCreateContext(NULL, 1, &clDevice, NULL, NULL, &errcode); CHECKERR(errcode); clCommands = clCreateCommandQueue(clContext, clDevice, CL_QUEUE_PROFILING_ENABLE, &errcode); CHECKERR(errcode); kernelFile = fopen("lud_kernel.cl", "r"); fseek(kernelFile, 0, SEEK_END); kernelLength = (size_t) ftell(kernelFile); kernelSource = (char *) malloc(sizeof(char)*kernelLength); rewind(kernelFile); fread((void *) kernelSource, kernelLength, 1, kernelFile); fclose(kernelFile); clProgram = clCreateProgramWithSource(clContext, 1, (const char **) &kernelSource, &kernelLength, &errcode); CHECKERR(errcode); free(kernelSource); char arg[100]; sprintf(arg,"-D BLOCK_SIZE=%d", (int)BLOCK_SIZE); errcode = clBuildProgram(clProgram, 1, &clDevice, arg, NULL, NULL); if (errcode == CL_BUILD_PROGRAM_FAILURE) { char *log; size_t logLength; errcode = clGetProgramBuildInfo(clProgram, clDevice, CL_PROGRAM_BUILD_LOG, 0, NULL, &logLength); log = (char *) malloc(sizeof(char)*logLength); errcode = clGetProgramBuildInfo(clProgram, clDevice, CL_PROGRAM_BUILD_LOG, logLength, (void *) log, NULL); fprintf(stderr, "Kernel build error! Log:\n%s", log); free(log); return 0; } CHECKERR(errcode); clKernel_diagonal = clCreateKernel(clProgram, "lud_diagonal", &errcode); CHECKERR(errcode); clKernel_perimeter = clCreateKernel(clProgram, "lud_perimeter", &errcode); CHECKERR(errcode); clKernel_internal = clCreateKernel(clProgram, "lud_internal", &errcode); CHECKERR(errcode); d_m = clCreateBuffer(clContext, CL_MEM_READ_WRITE, matrix_dim*matrix_dim*sizeof(float), NULL, &errcode); CHECKERR(errcode); /* beginning of timing point */ stopwatch_start(&sw); errcode = clEnqueueWriteBuffer(clCommands, d_m, CL_TRUE, 0, matrix_dim*matrix_dim*sizeof(float), (void *) m, 0, NULL, &ocdTempEvent); clFinish(clCommands); START_TIMER(ocdTempEvent, OCD_TIMER_H2D, "Matrix Copy", ocdTempTimer) END_TIMER(ocdTempTimer) CHECKERR(errcode); int i=0; size_t localWorkSize[2]; size_t globalWorkSize[2]; //printf("BLOCK_SIZE: %d\n",BLOCK_SIZE); // printf("max Work-item Size: %d\n",(int)max_worksize[0]); #ifdef START_POWER for( int iter = 0; iter < 1000; iter++) #endif for (i=0; i < matrix_dim-BLOCK_SIZE; i += BLOCK_SIZE) { errcode = clSetKernelArg(clKernel_diagonal, 0, sizeof(cl_mem), (void *) &d_m); errcode |= clSetKernelArg(clKernel_diagonal, 1, sizeof(int), (void *) &matrix_dim); errcode |= clSetKernelArg(clKernel_diagonal, 2, sizeof(int), (void *) &i); CHECKERR(errcode); localWorkSize[0] = BLOCK_SIZE; globalWorkSize[0] = BLOCK_SIZE; errcode = clEnqueueNDRangeKernel(clCommands, clKernel_diagonal, 1, NULL, globalWorkSize, localWorkSize, 0, NULL, &ocdTempEvent); clFinish(clCommands); START_TIMER(ocdTempEvent, OCD_TIMER_KERNEL, "Diagonal Kernels", ocdTempTimer) END_TIMER(ocdTempTimer) CHECKERR(errcode); errcode = clSetKernelArg(clKernel_perimeter, 0, sizeof(cl_mem), (void *) &d_m); errcode |= clSetKernelArg(clKernel_perimeter, 1, sizeof(int), (void *) &matrix_dim); errcode |= clSetKernelArg(clKernel_perimeter, 2, sizeof(int), (void *) &i); CHECKERR(errcode); localWorkSize[0] = BLOCK_SIZE*2; globalWorkSize[0] = ((matrix_dim-i)/BLOCK_SIZE-1)*localWorkSize[0]; errcode = clEnqueueNDRangeKernel(clCommands, clKernel_perimeter, 1, NULL, globalWorkSize, localWorkSize, 0, NULL, &ocdTempEvent); clFinish(clCommands); START_TIMER(ocdTempEvent, OCD_TIMER_KERNEL, "Perimeter Kernel", ocdTempTimer) CHECKERR(errcode); END_TIMER(ocdTempTimer) errcode = clSetKernelArg(clKernel_internal, 0, sizeof(cl_mem), (void *) &d_m); errcode |= clSetKernelArg(clKernel_internal, 1, sizeof(int), (void *) &matrix_dim); errcode |= clSetKernelArg(clKernel_internal, 2, sizeof(int), (void *) &i); CHECKERR(errcode); localWorkSize[0] = BLOCK_SIZE; localWorkSize[1] = BLOCK_SIZE; globalWorkSize[0] = ((matrix_dim-i)/BLOCK_SIZE-1)*localWorkSize[0]; globalWorkSize[1] = ((matrix_dim-i)/BLOCK_SIZE-1)*localWorkSize[1]; errcode = clEnqueueNDRangeKernel(clCommands, clKernel_internal, 2, NULL, globalWorkSize, localWorkSize, 0, NULL, &ocdTempEvent); clFinish(clCommands); START_TIMER(ocdTempEvent, OCD_TIMER_KERNEL, "Internal Kernel", ocdTempTimer) END_TIMER(ocdTempTimer) CHECKERR(errcode); } errcode = clSetKernelArg(clKernel_diagonal, 0, sizeof(cl_mem), (void *) &d_m); errcode |= clSetKernelArg(clKernel_diagonal, 1, sizeof(int), (void *) &matrix_dim); errcode |= clSetKernelArg(clKernel_diagonal, 2, sizeof(int), (void *) &i); CHECKERR(errcode); localWorkSize[0] = BLOCK_SIZE; globalWorkSize[0] = BLOCK_SIZE; errcode = clEnqueueNDRangeKernel(clCommands, clKernel_diagonal, 1, NULL, globalWorkSize, localWorkSize, 0, NULL, &ocdTempEvent); clFinish(clCommands); START_TIMER(ocdTempEvent, OCD_TIMER_KERNEL, "Diagonal Kernels", ocdTempTimer) CHECKERR(errcode); END_TIMER(ocdTempTimer) errcode = clEnqueueReadBuffer(clCommands, d_m, CL_TRUE, 0, matrix_dim*matrix_dim*sizeof(float), (void *) m, 0, NULL, &ocdTempEvent); clFinish(clCommands); START_TIMER(ocdTempEvent, OCD_TIMER_D2H, "Matrix copy", ocdTempTimer) END_TIMER(ocdTempTimer) /* end of timing point */ stopwatch_stop(&sw); printf("Time consumed(ms): %lf\n", 1000*get_interval_by_sec(&sw)); clReleaseMemObject(d_m); if (do_verify) { printf("After LUD\n"); print_matrix(m, matrix_dim); printf(">>>Verify<<<<\n"); printf("matrix_dim: %d\n",matrix_dim); lud_verify(mm, m, matrix_dim); free(mm); } clReleaseKernel(clKernel_diagonal); clReleaseKernel(clKernel_perimeter); clReleaseKernel(clKernel_internal); clReleaseProgram(clProgram); clReleaseCommandQueue(clCommands); clReleaseContext(clContext); free(m); ocd_finalize(); return EXIT_SUCCESS; } /* ---------- end of function main ---------- */
int main(int argc, char** argv) { ocd_init(&argc, &argv, NULL); ocd_initCL(); std::cerr << "N-Queen solver for OpenCL\n"; std::cerr << "Ping-Che Chen\n\n"; if(argc < 2) { std::cerr << "Usage: " << argv[0] << " [options] N\n"; std::cerr << "\tN: board size (1 ~ 32)\n"; std::cerr << "\t-cpu: use CPU (multi-threaded on Windows)\n"; std::cerr << "\t-prof: enable profiler\n"; std::cerr << "\t-threads #: set number of threads to #\n"; std::cerr << "\t-blocksize #: set size of thread blocks to #\n"; std::cerr << "\t-local: use local memory for arrays (default: off)\n"; std::cerr << "\t-noatomics: do not use global atomics\n"; std::cerr << "\t-novec: do not use vectorization\n"; std::cerr << "\t-vec4: use 4D vectors instead of 2D (only when vectorized- default: off)\n"; return 0; } // handle options bool force_cpu = false; bool profiling = false; int threads = 0; int block_size = 0; bool local = false;//default OFF (was true) bool noatomics = false; bool novec = false; bool use_vec4 = false; int start = 1; while(start < argc - 1) { if(std::strcmp(argv[start], "-cpu") == 0) { force_cpu = true; } else if(std::strcmp(argv[start], "-threads") == 0 && start < argc - 2) { threads = std::atoi(argv[start + 1]); start++; } else if(std::strcmp(argv[start], "-blocksize") == 0 && start < argc - 2) { block_size = std::atoi(argv[start + 1]); start++; } else if(std::strcmp(argv[start], "-local") == 0) { local = true; } else if(std::strcmp(argv[start], "-noatomics") == 0) { noatomics = true; } else if(std::strcmp(argv[start], "-novec") == 0) { novec = true; } else if(std::strcmp(argv[start], "-vec4") == 0) { use_vec4 = true; } else { std::cerr << "Unknown option " << argv[start] << "\n"; } start ++; } int board_size = std::atoi(argv[start]); if(board_size < 1 || board_size > 32) { std::cerr << "Inalid board size (only 1 ~ 32 allowed)\n"; return 0; } stopwatch sw; long long solutions = 0; long long unique_solutions = 0; if(force_cpu) { stopwatch_start(&sw); solutions = nqueen_cpu(board_size, &unique_solutions); stopwatch_stop(&sw); } else { stopwatch_start(&sw); cl_int err; // show device list size_t num_devices; num_devices=1;//In OpenDwarfs we only work with one device at a time. std::vector<cl_device_id> devices(num_devices / sizeof(cl_device_id)); devices.clear(); devices.resize(1); devices[0] = device_id; try { NQueenSolver nqueen(context, devices, profiling, threads, block_size, local, noatomics, novec, use_vec4); for(int i = 0; i < devices.size(); i++) { size_t name_length; err = clGetDeviceInfo(devices[i], CL_DEVICE_NAME, 0, 0, &name_length); if(err == CL_SUCCESS) { std::string name; name.resize(name_length + 1); clGetDeviceInfo(devices[i], CL_DEVICE_NAME, name_length, &name[0], &name_length); name[name_length] = 0; std::cerr << "Device " << i << ": " << name.c_str() << "\n"; std::cerr << "\tUsing " << nqueen.GetThreads(i) << " threads\n"; std::cerr << "\tBlock size = " << nqueen.GetBlockSize(i) << " threads\n"; if(nqueen.AtomicsEnabled(i)) { std::cerr << "\tUsing global atomics\n"; } if(nqueen.VectorizationEnabled(i)) { std::cerr << "\tUsing vectorization\n"; if(use_vec4) { std::cerr << "\tUse 4D vectors\n"; } else { std::cerr << "\tUse 2D vectors\n"; } } } } //start_time = std::clock(); solutions = nqueen.Compute(board_size, &unique_solutions); //end_time = std::clock(); } catch(CLError x) { if(x.GetErrorNo() == 1) { std::cerr << "1 OpenCL kernel execution failed\n"; } if(x.GetErrorNo() == 2) { std::cerr << "2 OpenCL kernel execution failed\n"; } if(x.GetErrorNo() == 3) { std::cerr << "3 OpenCL kernel execution failed\n"; } else { std::cerr << x << "\n"; } } stopwatch_stop(&sw); clReleaseContext(context); } std::cerr << "Solution took " << get_interval_by_sec(&sw) << " seconds to complete\n"; std::cerr << board_size << "-queen has " << solutions << " solutions (" << unique_solutions << " unique)\n"; printf("{ \"status\": %d, \"options\": \"-s %d\", \"time\": %f }\n", 1, board_size, get_interval_by_sec(&sw)); ocd_finalize(); return 0; }
int main ( int argc, char *argv[] ) { int matrix_dim = 32; /* default matrix_dim */ int opt, option_index=0, error=0; func_ret_t ret; const char *input_file = NULL; double *m, *mm; stopwatch sw; int i; while ((opt = getopt_long(argc, argv, ":vs:i:", long_options, &option_index)) != -1 ) { switch(opt){ case 'v': do_verify = 1; break; case 's': matrix_dim = atoi(optarg); break; case '?': fprintf(stderr, "invalid option\n"); error=1; break; case ':': fprintf(stderr, "missing argument\n"); error=1; break; default: error=1; } } if ((optind < argc) || (optind == 1) || error) { fprintf(stderr, "Usage: %s [-v] [-s matrix_size]\n", argv[0]); exit(EXIT_FAILURE); } if(matrix_dim>1) { fprintf(stderr, "Generating matrix of size %d x %d\n", matrix_dim, matrix_dim); ret = create_matrix_from_random(&m, matrix_dim); if(ret != RET_SUCCESS){ m = NULL; fprintf(stderr, "error could not generate random matrix of size %d x %d!\n", matrix_dim, matrix_dim); exit(EXIT_FAILURE); } } else { fprintf(stderr, "No input file or valid matrix size specified!\n"); exit(EXIT_FAILURE); } if (do_verify){ //printf("Before LUD\n"); //print_matrix(m, matrix_dim); matrix_duplicate(m, &mm, matrix_dim); } stopwatch_start(&sw); lud_base(m, matrix_dim); stopwatch_stop(&sw); if (matrix_dim == 1024) { for (i=0; i<100; ++i) { if (m[expected_row_indices[i]*matrix_dim + expected_col_indices[i]] != expected_values[i]) { fprintf(stderr, "ERROR: value at index (%d,%d) = '%.*f' is different from the expected value '%.*f'\n", expected_row_indices[i], expected_col_indices[i], // the 21 parameter prints enough significant decimal digits to obtain the same floating-point number // when read back 21, m[expected_row_indices[i]*matrix_dim + expected_col_indices[i]], 21, expected_values[i] ); fprintf(stderr, "Received values:\n"); for (i=0; i<100; ++i) { fprintf(stderr, "%.*f, ", 21, m[expected_row_indices[i]*matrix_dim + expected_col_indices[i]]); } fprintf(stderr, "\n"); exit(1); } } } else { fprintf(stderr, "WARNING: No self-checking step for dimension '%d'\n", matrix_dim); } if (do_verify){ //fprintf(stderr, "After LUD\n"); //print_matrix(m, matrix_dim); fprintf(stderr, ">>>Verify<<<<\n"); lud_verify(mm, m, matrix_dim); free(mm); } free(m); printf("{ \"status\": %d, \"options\": \"-s %d\", \"time\": %f }\n", 1, matrix_dim, get_interval_by_sec(&sw)); return EXIT_SUCCESS; } /* ---------- end of function main ---------- */
int main ( int argc, char *argv[] ) { printf("WG size of kernel = %d X %d\n", BLOCK_SIZE, BLOCK_SIZE); int matrix_dim = 32; /* default matrix_dim */ int opt, option_index=0; func_ret_t ret; const char *input_file = NULL; float *m, *mm; stopwatch sw; while ((opt = getopt_long(argc, argv, "::vs:i:", long_options, &option_index)) != -1 ) { switch(opt){ case 'i': input_file = optarg; break; case 'v': do_verify = 1; break; case 's': matrix_dim = atoi(optarg); printf("Generate input matrix internally, size =%d\n", matrix_dim); // fprintf(stderr, "Currently not supported, use -i instead\n"); // fprintf(stderr, "Usage: %s [-v] [-s matrix_size|-i input_file]\n", argv[0]); // exit(EXIT_FAILURE); break; case '?': fprintf(stderr, "invalid option\n"); break; case ':': fprintf(stderr, "missing argument\n"); break; default: fprintf(stderr, "Usage: %s [-v] [-s matrix_size|-i input_file]\n", argv[0]); exit(EXIT_FAILURE); } } if ( (optind < argc) || (optind == 1)) { fprintf(stderr, "Usage: %s [-v] [-s matrix_size|-i input_file]\n", argv[0]); exit(EXIT_FAILURE); } if (input_file) { printf("Reading matrix from file %s\n", input_file); ret = create_matrix_from_file(&m, input_file, &matrix_dim); if (ret != RET_SUCCESS) { m = NULL; fprintf(stderr, "error create matrix from file %s\n", input_file); exit(EXIT_FAILURE); } } else if (matrix_dim) { printf("Creating matrix internally size=%d\n", matrix_dim); ret = create_matrix(&m, matrix_dim); if (ret != RET_SUCCESS) { m = NULL; fprintf(stderr, "error create matrix internally size=%d\n", matrix_dim); exit(EXIT_FAILURE); } } else { printf("No input file specified!\n"); exit(EXIT_FAILURE); } if (do_verify){ printf("Before LUD\n"); // print_matrix(m, matrix_dim); matrix_duplicate(m, &mm, matrix_dim); } int sourcesize = 1024*1024; char * source = (char *)calloc(sourcesize, sizeof(char)); if(!source) { printf("ERROR: calloc(%d) failed\n", sourcesize); return -1; } char * kernel_lud_diag = "lud_diagonal"; char * kernel_lud_peri = "lud_perimeter"; char * kernel_lud_inter = "lud_internal"; FILE * fp = fopen("./lud_kernel.cl", "rb"); if(!fp) { printf("ERROR: unable to open '%s'\n"); return -1; } fread(source + strlen(source), sourcesize, 1, fp); fclose(fp); // Use 1: GPU 0: CPU int use_gpu = 1; // OpenCL initialization if(initialize(use_gpu)) return -1; // compile kernel cl_int err = 0; const char * slist[2] = { source, 0 }; cl_program prog = clCreateProgramWithSource(context, 1, slist, NULL, &err); if(err != CL_SUCCESS) { printf("ERROR: clCreateProgramWithSource() => %d\n", err); return -1; } char clOptions[110]; // sprintf(clOptions,"-I../../src"); sprintf(clOptions," "); #ifdef BLOCK_SIZE sprintf(clOptions + strlen(clOptions), " -DBLOCK_SIZE=%d", BLOCK_SIZE); #endif err = clBuildProgram(prog, 0, NULL, clOptions, NULL, NULL); { // show warnings/errors //static char log[65536]; memset(log, 0, sizeof(log)); //cl_device_id device_id = 0; //err = clGetContextInfo(context, CL_CONTEXT_DEVICES, sizeof(device_id), &device_id, NULL); //clGetProgramBuildInfo(prog, device_id, CL_PROGRAM_BUILD_LOG, sizeof(log)-1, log, NULL); //if(err || strstr(log,"warning:") || strstr(log, "error:")) printf("<<<<\n%s\n>>>>\n", log); } if(err != CL_SUCCESS) { printf("ERROR: clBuildProgram() => %d\n", err); return -1; } cl_kernel diagnal; cl_kernel perimeter; cl_kernel internal; diagnal = clCreateKernel(prog, kernel_lud_diag, &err); perimeter = clCreateKernel(prog, kernel_lud_peri, &err); internal = clCreateKernel(prog, kernel_lud_inter, &err); if(err != CL_SUCCESS) { printf("ERROR: clCreateKernel() 0 => %d\n", err); return -1; } clReleaseProgram(prog); //size_t local_work[3] = { 1, 1, 1 }; //size_t global_work[3] = {1, 1, 1 }; cl_mem d_m; d_m = clCreateBuffer(context, CL_MEM_READ_WRITE, matrix_dim*matrix_dim * sizeof(float), NULL, &err ); if(err != CL_SUCCESS) { printf("ERROR: clCreateBuffer d_m (size:%d) => %d\n", matrix_dim*matrix_dim, err); return -1;} /* beginning of timing point */ stopwatch_start(&sw); err = clEnqueueWriteBuffer(cmd_queue, d_m, 1, 0, matrix_dim*matrix_dim*sizeof(float), m, 0, 0, 0); if(err != CL_SUCCESS) { printf("ERROR: clEnqueueWriteBuffer d_m (size:%d) => %d\n", matrix_dim*matrix_dim, err); return -1; } int i=0; for (i=0; i < matrix_dim-BLOCK_SIZE; i += BLOCK_SIZE) { clSetKernelArg(diagnal, 0, sizeof(void *), (void*) &d_m); clSetKernelArg(diagnal, 1, sizeof(float) * BLOCK_SIZE * BLOCK_SIZE, (void*)NULL ); clSetKernelArg(diagnal, 2, sizeof(cl_int), (void*) &matrix_dim); clSetKernelArg(diagnal, 3, sizeof(cl_int), (void*) &i); size_t global_work1[3] = {BLOCK_SIZE, 1, 1}; size_t local_work1[3] = {BLOCK_SIZE, 1, 1}; err = clEnqueueNDRangeKernel(cmd_queue, diagnal, 2, NULL, global_work1, local_work1, 0, 0, 0); if(err != CL_SUCCESS) { printf("ERROR: diagnal clEnqueueNDRangeKernel()=>%d failed\n", err); return -1; } clSetKernelArg(perimeter, 0, sizeof(void *), (void*) &d_m); clSetKernelArg(perimeter, 1, sizeof(float) * BLOCK_SIZE * BLOCK_SIZE, (void*)NULL ); clSetKernelArg(perimeter, 2, sizeof(float) * BLOCK_SIZE * BLOCK_SIZE, (void*)NULL ); clSetKernelArg(perimeter, 3, sizeof(float) * BLOCK_SIZE * BLOCK_SIZE, (void*)NULL ); clSetKernelArg(perimeter, 4, sizeof(cl_int), (void*) &matrix_dim); clSetKernelArg(perimeter, 5, sizeof(cl_int), (void*) &i); size_t global_work2[3] = {BLOCK_SIZE * 2 * ((matrix_dim-i)/BLOCK_SIZE-1), 1, 1}; size_t local_work2[3] = {BLOCK_SIZE * 2, 1, 1}; err = clEnqueueNDRangeKernel(cmd_queue, perimeter, 2, NULL, global_work2, local_work2, 0, 0, 0); if(err != CL_SUCCESS) { printf("ERROR: perimeter clEnqueueNDRangeKernel()=>%d failed\n", err); return -1; } clSetKernelArg(internal, 0, sizeof(void *), (void*) &d_m); clSetKernelArg(internal, 1, sizeof(float) * BLOCK_SIZE * BLOCK_SIZE, (void*)NULL ); clSetKernelArg(internal, 2, sizeof(float) * BLOCK_SIZE * BLOCK_SIZE, (void*)NULL ); clSetKernelArg(internal, 3, sizeof(cl_int), (void*) &matrix_dim); clSetKernelArg(internal, 4, sizeof(cl_int), (void*) &i); size_t global_work3[3] = {BLOCK_SIZE * ((matrix_dim-i)/BLOCK_SIZE-1), BLOCK_SIZE * ((matrix_dim-i)/BLOCK_SIZE-1), 1}; size_t local_work3[3] = {BLOCK_SIZE, BLOCK_SIZE, 1}; err = clEnqueueNDRangeKernel(cmd_queue, internal, 2, NULL, global_work3, local_work3, 0, 0, 0); if(err != CL_SUCCESS) { printf("ERROR: internal clEnqueueNDRangeKernel()=>%d failed\n", err); return -1; } } clSetKernelArg(diagnal, 0, sizeof(void *), (void*) &d_m); clSetKernelArg(diagnal, 1, sizeof(float) * BLOCK_SIZE * BLOCK_SIZE, (void*)NULL ); clSetKernelArg(diagnal, 2, sizeof(cl_int), (void*) &matrix_dim); clSetKernelArg(diagnal, 3, sizeof(cl_int), (void*) &i); size_t global_work1[3] = {BLOCK_SIZE, 1, 1}; size_t local_work1[3] = {BLOCK_SIZE, 1, 1}; err = clEnqueueNDRangeKernel(cmd_queue, diagnal, 2, NULL, global_work1, local_work1, 0, 0, 0); if(err != CL_SUCCESS) { printf("ERROR: diagnal clEnqueueNDRangeKernel()=>%d failed\n", err); return -1; } err = clEnqueueReadBuffer(cmd_queue, d_m, 1, 0, matrix_dim*matrix_dim*sizeof(float), m, 0, 0, 0); if(err != CL_SUCCESS) { printf("ERROR: clEnqueueReadBuffer d_m (size:%d) => %d\n", matrix_dim*matrix_dim, err); return -1; } clFinish(cmd_queue); /* end of timing point */ stopwatch_stop(&sw); printf("Time consumed(ms): %lf\n", 1000*get_interval_by_sec(&sw)); clReleaseMemObject(d_m); if (do_verify){ printf("After LUD\n"); // print_matrix(m, matrix_dim); printf(">>>Verify<<<<\n"); lud_verify(mm, m, matrix_dim); free(mm); } free(m); if(shutdown()) return -1; }
int main ( int argc, char *argv[] ) { //printf("Starting..\n"); int matrix_dim = 32; /* default size */ int opt, option_index=0; func_ret_t ret; const char *input_file = NULL; float *m, *mm; stopwatch sw; int grid_x=0; int grid_y=0; while ((opt = getopt_long(argc, argv, "::vs:i:x:y:", long_options, &option_index)) != -1 ) { switch(opt){ case 'i': input_file = optarg; break; case 'v': do_verify = 1; break; case 's': matrix_dim = atoi(optarg); //printf("Generate input matrix internally, size =%d\n", matrix_dim); // fprintf(stderr, "Currently not supported, use -i instead\n"); // fprintf(stderr, "Usage: %s [-v] [-s matrix_size|-i input_file]\n", argv[0]); // exit(EXIT_FAILURE); break; case 'x': grid_x = atoi(optarg); break; case 'y': grid_y = atoi(optarg); break; case '?': fprintf(stderr, "invalid option\n"); break; case ':': fprintf(stderr, "missing argument\n"); break; default: fprintf(stderr, "1Usage: %s [-v] [-s matrix_size|-i input_file]\n", argv[0]); exit(EXIT_FAILURE); } } /* if ( (optind < argc) || (optind == 1)) { fprintf(stderr, "2Usage: %s [-v] [-n no. of threads] [-s matrix_size|-i input_file]\n", argv[0]); exit(EXIT_FAILURE); } */ if (input_file) { //printf("Reading matrix from file %s\n", input_file); ret = create_matrix_from_file(&m, input_file, &matrix_dim); if (ret != RET_SUCCESS) { m = NULL; fprintf(stderr, "error create matrix from file %s\n", input_file); exit(EXIT_FAILURE); } } else if (matrix_dim) { //printf("Creating matrix internally size=%d\n", matrix_dim); ret = create_matrix(&m, matrix_dim); if (ret != RET_SUCCESS) { m = NULL; fprintf(stderr, "error create matrix internally size=%d\n", matrix_dim); exit(EXIT_FAILURE); } } else { printf("No input file specified!\n"); exit(EXIT_FAILURE); } if (do_verify){ /* print_matrix(m, matrix_dim); */ matrix_duplicate(m, &mm, matrix_dim); } wul(); //printf("Starting. . . \n"); //lud_oacc(m, matrix_dim,grid_x,grid_y); stopwatch_start(&sw); // lud_omp(m, matrix_dim); lud_oacc(m, matrix_dim,grid_x,grid_y); stopwatch_stop(&sw); printf("Time consumed(ms): %lf\n", 1000*get_interval_by_sec(&sw)); if (do_verify){ printf("After LUD\n"); /* print_matrix(m, matrix_dim); */ printf(">>>Verify<<<<\n"); lud_verify(mm, m, matrix_dim); free(mm); } free(m); return EXIT_SUCCESS; } /* ---------- end of function main ---------- */
int main ( int argc, char *argv[] ) { int matrix_dim = 32; /* default matrix_dim */ int opt, option_index=0, error=0; func_ret_t ret; char *input_path = NULL; char *output_path = NULL; FILE *file; double *m; stopwatch sw; int i,j; int debug = 0; size_t linesiz=0; char* linebuf=NULL; ssize_t linelen=0; char* token; const char comma[2] = ","; while ((opt = getopt_long(argc, argv, ":dvs:i:o:", long_options, &option_index)) != -1 ) { switch(opt){ case 'v': do_verify = 1; break; case 's': matrix_dim = atoi(optarg); break; case 'i': input_path = optarg; break; case 'o': output_path = optarg; break; case '?': fprintf(stderr, "invalid option\n"); error=1; break; case 'd': debug=1; break; case ':': fprintf(stderr, "missing argument\n"); error=1; break; default: error=1; } } if ((optind < argc) || (optind == 1) || input_path == NULL || error) { fprintf(stderr, "Usage: %s -s size -i input_path [-v] [-d] [-o output_path]\n", argv[0]); exit(EXIT_FAILURE); } file = fopen(input_path, "r"); if (file == NULL) { fprintf(stderr, "Invalid input file path: %s\n", input_path); exit(EXIT_FAILURE); } // Read matrix from file if (debug) { fprintf(stderr, "Reading data from file %s\n", input_path); } m = (double *)malloc(sizeof(double) * matrix_dim * matrix_dim); j = 0; while ((linelen=getline(&linebuf, &linesiz, file)) != -1) { if (debug) { fprintf(stderr, "Read line: %s\n", linebuf); } /* get the first number */ i = 0; token = strtok(linebuf, comma); /* walk through other numbers */ while( token != NULL ) { if (debug) { fprintf(stderr, "Read token: %s\n", token); } m[j*matrix_dim + i] = atof(token); token = strtok(NULL, comma); i = i + 1; } free(linebuf); linebuf=NULL; j = j + 1; } fclose(file); if (debug) { fprintf(stderr, "Computing LUD\n"); } stopwatch_start(&sw); lud_base(m, matrix_dim); stopwatch_stop(&sw); if (output_path) { if (debug) { fprintf(stderr, "Saving result in %s\n", output_path); } file = fopen(output_path, "w"); if (file == NULL) { free(m); exit(EXIT_FAILURE); } for (j = 0; j < matrix_dim; ++j) { for (i = 0; i < matrix_dim; ++i) { fprintf(file, "%.*f", 21, m[j*matrix_dim+i]); if (i < matrix_dim-1) { fprintf(file, ","); } } fprintf(file, "\n"); } fclose(file); } free(m); printf("{ \"status\": %d, \"options\": \"-s %d\", \"time\": %f }\n", 1, matrix_dim, get_interval_by_sec(&sw)); return EXIT_SUCCESS; } /* ---------- end of function main ---------- */