diff --git a/XSbench_shared_header.h b/XSbench_shared_header.h index 779e8f39..89ba59ca 100644 --- a/XSbench_shared_header.h +++ b/XSbench_shared_header.h @@ -3,23 +3,6 @@ // Header for shared utilities across XSBench versions -typedef struct{ - int nthreads; - long n_isotopes; - long n_gridpoints; - int lookups; - char * HM; - int grid_type; // 0: Unionized Grid (default) 1: Nuclide Grid - int hash_bins; - int particles; - int simulation_method; - int binary_mode; - int kernel_id; - int num_iterations; - int num_warmups; - char *filename; -} Inputs; - typedef struct{ double device_to_host_time; double kernel_time; diff --git a/cuda/Main.cu b/cuda/Main.cu index c8c29af2..c1f8ef28 100644 --- a/cuda/Main.cu +++ b/cuda/Main.cu @@ -37,8 +37,6 @@ int main(int argc, char *argv[]) { if (in.binary_mode == WRITE && mype == 0) binary_write(in, SD); - Profile profile; - // ===================================================================== // Cross Section (XS) Parallel Lookup Simulation // This is the section that should be profiled, as it reflects a @@ -58,7 +56,7 @@ int main(int argc, char *argv[]) { // Run simulation if (in.simulation_method == EVENT_BASED) { if (in.kernel_id == 0) - verification = run_event_based_simulation_baseline(in, SD, mype, &profile); + verification = run_event_based_simulation_baseline(in, SD, mype); else if (in.kernel_id == 1) verification = run_event_based_simulation_optimization_1(in, SD, mype); else if (in.kernel_id == 2) @@ -101,7 +99,5 @@ int main(int argc, char *argv[]) { int is_invalid_result = print_results(in, mype, omp_end - omp_start, nprocs, verification); - print_profile(profile, in); - return is_invalid_result; } diff --git a/cuda/Simulation.cu b/cuda/Simulation.cu index ed2fcee7..2070a751 100644 --- a/cuda/Simulation.cu +++ b/cuda/Simulation.cu @@ -12,12 +12,10 @@ // line argument. //////////////////////////////////////////////////////////////////////////////////// -unsigned long long run_event_based_simulation_baseline(Inputs in, SimulationData SD, int mype, Profile* profile) +unsigned long long run_event_based_simulation_baseline(Inputs in, SimulationData SD, int mype) { - double start = get_time(); // Move Data to GPU SimulationData GSD = move_simulation_data_to_device(in, mype, SD); - profile->host_to_device_time = get_time() - start; //////////////////////////////////////////////////////////////////////////////// // Configure & Launch Simulation Kernel @@ -27,27 +25,17 @@ unsigned long long run_event_based_simulation_baseline(Inputs in, SimulationData int nthreads = 256; int nblocks = ceil( (double) in.lookups / (double) nthreads); - int nwarmups = in.num_warmups; start = 0.0; - for (int i = 0; i < in.num_iterations + nwarmups; i++) { - if (i == nwarmups) { - gpuErrchk( cudaDeviceSynchronize() ); - start = get_time(); - } - xs_lookup_kernel_baseline<<>>( in, GSD ); - } + xs_lookup_kernel_baseline<<>>( in, GSD ); gpuErrchk( cudaPeekAtLastError() ); gpuErrchk( cudaDeviceSynchronize() ); - profile->kernel_time = get_time() - start; //////////////////////////////////////////////////////////////////////////////// // Reduce Verification Results //////////////////////////////////////////////////////////////////////////////// if( mype == 0) printf("Reducing verification results...\n"); - start = get_time(); gpuErrchk(cudaMemcpy(SD.verification, GSD.verification, in.lookups * sizeof(unsigned long), cudaMemcpyDeviceToHost) ); - profile->device_to_host_time = get_time() - start; unsigned long verification_scalar = 0; for( int i =0; i < in.lookups; i++ ) diff --git a/cuda/XSbench_header.cuh b/cuda/XSbench_header.cuh index 2d678777..25e5fec7 100644 --- a/cuda/XSbench_header.cuh +++ b/cuda/XSbench_header.cuh @@ -10,7 +10,6 @@ #include #include #include -#include "../XSbench_shared_header.h" // Grid types #define UNIONIZED 0 @@ -49,6 +48,20 @@ typedef struct{ double nu_fission_xs; } NuclideGridPoint; +typedef struct{ + int nthreads; + long n_isotopes; + long n_gridpoints; + int lookups; + char * HM; + int grid_type; // 0: Unionized Grid (default) 1: Nuclide Grid + int hash_bins; + int particles; + int simulation_method; + int binary_mode; + int kernel_id; +} Inputs; + typedef struct{ int * num_nucs; // Length = length_num_nucs; double * concs; // Length = length_concs @@ -84,7 +97,7 @@ void binary_write( Inputs in, SimulationData SD ); SimulationData binary_read( Inputs in ); // Simulation.cu -unsigned long long run_event_based_simulation_baseline(Inputs in, SimulationData SD, int mype, Profile* profile); +unsigned long long run_event_based_simulation_baseline(Inputs in, SimulationData SD, int mype); __global__ void xs_lookup_kernel_baseline(Inputs in, SimulationData GSD ); __device__ void calculate_micro_xs( double p_energy, int nuc, long n_isotopes, long n_gridpoints, diff --git a/cuda/io.cu b/cuda/io.cu index 320fde60..5d7dcea0 100644 --- a/cuda/io.cu +++ b/cuda/io.cu @@ -158,7 +158,6 @@ void print_inputs(Inputs in, int nprocs, int version ) printf("XS Lookups per Particle: "); fancy_int(in.lookups); } printf("Total XS Lookups: "); fancy_int(in.lookups); - printf("Total XS Iterations: "); fancy_int(in.num_iterations); #ifdef MPI printf("MPI Ranks: %d\n", nprocs); printf("Mem Usage per MPI Rank (MB): "); fancy_int(mem_tot); @@ -219,9 +218,6 @@ void print_CLI_error(void) printf(" -h Number of hash bins (only relevant when used with \"-G hash\")\n"); printf(" -b Read or write all data structures to file. If reading, this will skip initialization phase. (read, write)\n"); printf(" -k Specifies which kernel to run. 0 is baseline, 1, 2, etc are optimized variants. (0 is default.)\n"); - printf(" -n Specifies how many kernel iterations to run. (1 is default.)\n"); - printf(" -w Specifies how many warmup iterations to run. (0 is default.)\n"); - printf(" --csv Save output to csv file. (Default is stdout)\n"); printf("Default is equivalent to: -m history -s large -l 34 -p 500000 -G unionized -k 0 -n 1\n"); printf("See readme for full description of default run values\n"); exit(4); @@ -261,15 +257,6 @@ Inputs read_CLI( int argc, char * argv[] ) // defaults to baseline kernel input.kernel_id = 0; - // default to one kernel iteration - input.num_iterations = 1; - - // default to zero warmup iterations - input.num_warmups = 1; - - // default to stdout - input.filename = NULL; - // defaults to H-M Large benchmark input.HM = (char *) malloc( 6 * sizeof(char) ); input.HM[0] = 'l' ; @@ -407,36 +394,6 @@ Inputs read_CLI( int argc, char * argv[] ) else print_CLI_error(); } - // number of kernel iterations (-n) - else if( strcmp(arg, "-n") == 0 ) - { - if( ++i < argc) - { - input.num_iterations = atoi(argv[i]); - } - else - print_CLI_error(); - } - else if( strcmp(arg, "--csv") == 0 ) - { - if( ++i < argc ) { - input.filename = (char *)malloc(strlen(argv[i]) + 1); - strcpy(input.filename, argv[i]); - } - else - print_CLI_error(); - } - else if( strcmp(arg, "-w") == 0 ) - { - if( ++i < argc) - { - input.num_warmups = atoi(argv[i]); - } - else - print_CLI_error(); - } - else - print_CLI_error(); } // Validate Input @@ -461,10 +418,6 @@ Inputs read_CLI( int argc, char * argv[] ) if( input.hash_bins < 1 ) print_CLI_error(); - // Validate number of iterations - if ( input.num_iterations < 1 ) - print_CLI_error(); - // Validate HM size if( strcasecmp(input.HM, "small") != 0 && strcasecmp(input.HM, "large") != 0 && diff --git a/hip/Main.cpp b/hip/Main.cpp index 4c46699f..3b2264df 100644 --- a/hip/Main.cpp +++ b/hip/Main.cpp @@ -38,12 +38,8 @@ int main( int argc, char* argv[] ) if( in.binary_mode == WRITE && mype == 0 ) binary_write(in, SD); - Profile profile; - // Move data to GPU - double start = get_time(); SimulationData GSD = move_simulation_data_to_device( in, mype, SD ); - profile.host_to_device_time = get_time() - start; // ===================================================================== // Cross Section (XS) Parallel Lookup Simulation @@ -66,7 +62,7 @@ int main( int argc, char* argv[] ) if( in.simulation_method == EVENT_BASED ) { if( in.kernel_id == 0 ) - verification = run_event_based_simulation_baseline(in, GSD, mype, &profile); + verification = run_event_based_simulation_baseline(in, GSD, mype); /* else if( in.kernel_id == 1 ) verification = run_event_based_simulation_optimization_1(in, GSD, mype); @@ -109,7 +105,5 @@ int main( int argc, char* argv[] ) // Print / Save Results and Exit int is_invalid_result = print_results( in, mype, omp_end-omp_start, nprocs, verification ); - print_profile(profile, in); - return is_invalid_result; } diff --git a/hip/Simulation.cpp b/hip/Simulation.cpp index ae601037..08575a1c 100644 --- a/hip/Simulation.cpp +++ b/hip/Simulation.cpp @@ -13,7 +13,7 @@ // line argument. //////////////////////////////////////////////////////////////////////////////////// -unsigned long long run_event_based_simulation_baseline(Inputs in, SimulationData GSD, int mype, Profile* profile) +unsigned long long run_event_based_simulation_baseline(Inputs in, SimulationData GSD, int mype) { //////////////////////////////////////////////////////////////////////////////// // Configure & Launch Simulation Kernel @@ -23,25 +23,14 @@ unsigned long long run_event_based_simulation_baseline(Inputs in, SimulationData int nthreads = 256; int nblocks = ceil( (double) in.lookups / (double) nthreads); - int nwarmups = in.num_warmups; - double start = 0.0; - for (int i = 0; i < in.num_iterations + nwarmups; i++) { - if (i == nwarmups) { - gpuErrchk( hipDeviceSynchronize() ); - start = get_time(); - } - hipLaunchKernelGGL(xs_lookup_kernel_baseline, dim3(nblocks), dim3(nthreads), 0, 0, in, GSD ); - } + hipLaunchKernelGGL(xs_lookup_kernel_baseline, dim3(nblocks), dim3(nthreads), 0, 0, in, GSD ); gpuErrchk( hipPeekAtLastError() ); gpuErrchk( hipDeviceSynchronize() ); - profile->kernel_time = get_time() - start; size_t sz = in.lookups * sizeof(unsigned long); unsigned long * v = (unsigned long *) malloc(sz); - start = get_time(); gpuErrchk( hipMemcpy(v, GSD.verification, sz, hipMemcpyDeviceToHost) ); - profile->device_to_host_time = get_time() - start; //////////////////////////////////////////////////////////////////////////////// // Reduce Verification Results diff --git a/hip/XSbench_header.h b/hip/XSbench_header.h index aea04d12..c69c59f0 100644 --- a/hip/XSbench_header.h +++ b/hip/XSbench_header.h @@ -9,7 +9,6 @@ #include #include #include -#include "../XSbench_shared_header.h" // Grid types #define UNIONIZED 0 @@ -48,6 +47,20 @@ typedef struct{ double nu_fission_xs; } NuclideGridPoint; +typedef struct{ + int nthreads; + long n_isotopes; + long n_gridpoints; + int lookups; + char * HM; + int grid_type; // 0: Unionized Grid (default) 1: Nuclide Grid + int hash_bins; + int particles; + int simulation_method; + int binary_mode; + int kernel_id; +} Inputs; + typedef struct{ int * num_nucs; // Length = length_num_nucs; double * concs; // Length = length_concs @@ -83,7 +96,7 @@ void binary_write( Inputs in, SimulationData SD ); SimulationData binary_read( Inputs in ); // Simulation.cu -unsigned long long run_event_based_simulation_baseline(Inputs in, SimulationData SD, int mype, Profile* profile); +unsigned long long run_event_based_simulation_baseline(Inputs in, SimulationData SD, int mype); __global__ void xs_lookup_kernel_baseline(Inputs in, SimulationData GSD ); __device__ void calculate_micro_xs( double p_energy, int nuc, long n_isotopes, long n_gridpoints, diff --git a/hip/io.cpp b/hip/io.cpp index 3afbf4b6..afb8f623 100644 --- a/hip/io.cpp +++ b/hip/io.cpp @@ -158,7 +158,6 @@ void print_inputs(Inputs in, int nprocs, int version ) printf("XS Lookups per Particle: "); fancy_int(in.lookups); } printf("Total XS Lookups: "); fancy_int(in.lookups); - printf("Total XS Iterations: "); fancy_int(in.num_iterations); #ifdef MPI printf("MPI Ranks: %d\n", nprocs); printf("Mem Usage per MPI Rank (MB): "); fancy_int(mem_tot); @@ -219,7 +218,6 @@ void print_CLI_error(void) printf(" -h Number of hash bins (only relevant when used with \"-G hash\")\n"); printf(" -b Read or write all data structures to file. If reading, this will skip initialization phase. (read, write)\n"); printf(" -k Specifies which kernel to run. 0 is baseline, 1, 2, etc are optimized variants. (0 is default.)\n"); - printf(" -n Specifies how many kernel iterations to run. (1 is default.)\n"); printf("Default is equivalent to: -m history -s large -l 34 -p 500000 -G unionized -k 0 -n 1\n"); printf("See readme for full description of default run values\n"); exit(4); @@ -259,15 +257,6 @@ Inputs read_CLI( int argc, char * argv[] ) // defaults to baseline kernel input.kernel_id = 0; - // defaults to one kernel iteration - input.num_iterations = 1; - - // default to zero warmup iterations - input.num_warmups = 1; - - // default to stdout - input.filename = NULL; - // defaults to H-M Large benchmark input.HM = (char *) malloc( 6 * sizeof(char) ); input.HM[0] = 'l' ; @@ -405,36 +394,6 @@ Inputs read_CLI( int argc, char * argv[] ) else print_CLI_error(); } - // number of kernel iterations (-n) - else if( strcmp(arg, "-n") == 0 ) - { - if( ++i < argc ) - { - input.num_iterations = atoi(argv[i]); - } - else - print_CLI_error(); - } - else if( strcmp(arg, "--csv") == 0 ) - { - if( ++i < argc ) { - input.filename = (char *)malloc(strlen(argv[i]) + 1); - strcpy(input.filename, argv[i]); - } - else - print_CLI_error(); - } - else if( strcmp(arg, "-w") == 0 ) - { - if( ++i < argc) - { - input.num_warmups = atoi(argv[i]); - } - else - print_CLI_error(); - } - else - print_CLI_error(); } // Validate Input @@ -459,10 +418,6 @@ Inputs read_CLI( int argc, char * argv[] ) if( input.hash_bins < 1 ) print_CLI_error(); - // Validate number of iterations - if( input.num_iterations < 1 ) - print_CLI_error(); - // Validate HM size if( strcasecmp(input.HM, "small") != 0 && strcasecmp(input.HM, "large") != 0 && diff --git a/kokkos/CMakeLists.txt b/kokkos/CMakeLists.txt index 450c17dc..8dd0e445 100644 --- a/kokkos/CMakeLists.txt +++ b/kokkos/CMakeLists.txt @@ -19,14 +19,14 @@ set(CMAKE_CXX_FLAGS_RELEASE "-O3") find_package(Kokkos REQUIRED) -if (${KOKKOS_BACK_END} STREQUAL "CUDA") +if (Kokkos_ENABLE_CUDA) enable_language(CUDA) set(CMAKE_CUDA_STANDARD 17) set(CMAKE_CUDA_FLAGS "${CMAKE_CUDA_FLAGS} -extended-lambda -Wext-lambda-captures-this -expt-relaxed-constexpr") set_source_files_properties(${SOURCE} PROPERTIES LANGUAGE CUDA) -elseif (${KOKKOS_BACK_END} STREQUAL "HIP") +elseif (Kokkos_ENABLE_HIP) enable_language(HIP) set(CMAKE_HIP_STANDARD 17) diff --git a/kokkos/Main.cpp b/kokkos/Main.cpp index eaedaf0d..3c42ed15 100644 --- a/kokkos/Main.cpp +++ b/kokkos/Main.cpp @@ -21,10 +21,8 @@ int main( int argc, char* argv[] ) MPI_Comm_rank(MPI_COMM_WORLD, &mype); #endif - Profile profile; - - // Start Kokkos - Kokkos::initialize(); + // Start Kokkos + Kokkos::initialize(); // Process CLI Fields -- store in "Inputs" structure Inputs in = read_CLI( argc, argv ); @@ -79,7 +77,7 @@ int main( int argc, char* argv[] ) if( in.simulation_method == EVENT_BASED ) { if( in.kernel_id == 0 ) - verification = run_event_based_simulation(in, SD, mype, &elapsed_time, &profile); + verification = run_event_based_simulation(in, SD, mype, &elapsed_time); else { printf("Error: No kernel ID %d found!\n", in.kernel_id); @@ -114,7 +112,5 @@ int main( int argc, char* argv[] ) MPI_Finalize(); #endif - print_profile(profile, in); - return is_invalid_result; } diff --git a/kokkos/Simulation.cpp b/kokkos/Simulation.cpp index 4ecf9663..29801ff4 100644 --- a/kokkos/Simulation.cpp +++ b/kokkos/Simulation.cpp @@ -11,7 +11,7 @@ // are not yet implemented for this OpenMP targeting offload port. //////////////////////////////////////////////////////////////////////////////////// -unsigned long long run_event_based_simulation(Inputs in, SimulationData SD, int mype, double* end, Profile* profile) +unsigned long long run_event_based_simulation(Inputs in, SimulationData SD, int mype, double* end) { if( mype == 0) printf("Beginning event based simulation...\n"); @@ -40,8 +40,6 @@ unsigned long long run_event_based_simulation(Inputs in, SimulationData SD, int Kokkos::Timer start; start.reset(); - double startP = get_time(); - UIntView u_num_nucs(SD.num_nucs, SD.length_num_nucs); SD.d_num_nucs = new IntView(Kokkos::ViewAllocateWithoutInitializing("d_num_nucs"), SD.length_num_nucs); @@ -92,81 +90,66 @@ unsigned long long run_event_based_simulation(Inputs in, SimulationData SD, int Kokkos::create_mirror_view(d_verification); Kokkos::deep_copy(d_verification, verification); - profile->host_to_device_time = get_time() - startP; - - int nwarmups = in.num_warmups; - for (int it = 0; it < in.num_iterations + nwarmups; it++) { - if (it == nwarmups) { - Kokkos::fence(); - startP = get_time(); - } - Kokkos::parallel_for("Simulation", - Kokkos::RangePolicy(0, in.lookups), - KOKKOS_LAMBDA (int i) { - // Set the initial seed value - uint64_t seed = STARTING_SEED; - - // Forward seed to lookup index (we need 2 samples per lookup) - seed = fast_forward_LCG(seed, 2*i); - - // Randomly pick an energy and material for the particle - double p_energy = LCG_random_double(&seed); - int mat = pick_mat(&seed); - - // debugging - //printf("E = %lf mat = %d\n", p_energy, mat); - - double macro_xs_vector[5] = {0}; - - // Perform macroscopic Cross Section Lookup - calculate_macro_xs( - p_energy, // Sampled neutron energy (in lethargy) - mat, // Sampled material type index neutron is in - in.n_isotopes, // Total number of isotopes in simulation - in.n_gridpoints, // Number of gridpoints per isotope in simulation - num_nucs, // 1-D array with number of nuclides per material - concs, // Flattened 2-D array with concentration of each nuclide in each material - unionized_energy_array, // 1-D Unionized energy array - index_grid, // Flattened 2-D grid holding indices into nuclide grid for each unionized energy level - nuclide_grid, // Flattened 2-D grid holding energy levels and XS_data for all nuclides in simulation - mats, // Flattened 2-D array with nuclide indices defining composition of each type of material - macro_xs_vector, // 1-D array with result of the macroscopic cross section (5 different reaction channels) - in.grid_type, // Lookup type (nuclide, hash, or unionized) - in.hash_bins, // Number of hash bins used (if using hash lookup type) - SD.max_num_nucs // Maximum number of nuclides present in any material - ); - - // For verification, and to prevent the compiler from optimizing - // all work out, we interrogate the returned macro_xs_vector array - // to find its maximum value index, then increment the verification - // value by that index. In this implementation, we prevent thread - // contention by using an OMP reduction on the verification value. - // For accelerators, a different approach might be required - // (e.g., atomics, reduction of thread-specific values in large - // array via CUDA thrust, etc). - double max = -1.0; - int max_idx = 0; - for(int j = 0; j < 5; j++ ) + Kokkos::parallel_for("Simulation", + Kokkos::RangePolicy(0, in.lookups), + KOKKOS_LAMBDA (int i) { + // Set the initial seed value + uint64_t seed = STARTING_SEED; + + // Forward seed to lookup index (we need 2 samples per lookup) + seed = fast_forward_LCG(seed, 2*i); + + // Randomly pick an energy and material for the particle + double p_energy = LCG_random_double(&seed); + int mat = pick_mat(&seed); + + // debugging + //printf("E = %lf mat = %d\n", p_energy, mat); + + double macro_xs_vector[5] = {0}; + + // Perform macroscopic Cross Section Lookup + calculate_macro_xs( + p_energy, // Sampled neutron energy (in lethargy) + mat, // Sampled material type index neutron is in + in.n_isotopes, // Total number of isotopes in simulation + in.n_gridpoints, // Number of gridpoints per isotope in simulation + num_nucs, // 1-D array with number of nuclides per material + concs, // Flattened 2-D array with concentration of each nuclide in each material + unionized_energy_array, // 1-D Unionized energy array + index_grid, // Flattened 2-D grid holding indices into nuclide grid for each unionized energy level + nuclide_grid, // Flattened 2-D grid holding energy levels and XS_data for all nuclides in simulation + mats, // Flattened 2-D array with nuclide indices defining composition of each type of material + macro_xs_vector, // 1-D array with result of the macroscopic cross section (5 different reaction channels) + in.grid_type, // Lookup type (nuclide, hash, or unionized) + in.hash_bins, // Number of hash bins used (if using hash lookup type) + SD.max_num_nucs // Maximum number of nuclides present in any material + ); + + // For verification, and to prevent the compiler from optimizing + // all work out, we interrogate the returned macro_xs_vector array + // to find its maximum value index, then increment the verification + // value by that index. In this implementation, we prevent thread + // contention by using an OMP reduction on the verification value. + // For accelerators, a different approach might be required + // (e.g., atomics, reduction of thread-specific values in large + // array via CUDA thrust, etc). + double max = -1.0; + int max_idx = 0; + for(int j = 0; j < 5; j++ ) + { + if( macro_xs_vector[j] > max ) { - if( macro_xs_vector[j] > max ) - { - max = macro_xs_vector[j]; - max_idx = j; - } + max = macro_xs_vector[j]; + max_idx = j; } - d_verification(i) = max_idx+1; - }); - } + } + d_verification(i) = max_idx+1; + }); Kokkos::fence(); - profile->kernel_time = get_time() - startP; - - startP = get_time(); - Kokkos::deep_copy(verification, d_verification); - profile->device_to_host_time = get_time() - startP; - // End Simulation Timer *end = start.seconds(); diff --git a/kokkos/XSbench_header.hpp b/kokkos/XSbench_header.hpp index 46444657..fc7e5de9 100644 --- a/kokkos/XSbench_header.hpp +++ b/kokkos/XSbench_header.hpp @@ -11,7 +11,6 @@ #include #include #include -#include "../XSbench_shared_header.h" #include @@ -47,6 +46,20 @@ typedef struct{ double nu_fission_xs; } NuclideGridPoint; +typedef struct{ + int nthreads; + long n_isotopes; + long n_gridpoints; + int lookups; + char * HM; + int grid_type; // 0: Unionized Grid (default) 1: Nuclide Grid + int hash_bins; + int particles; + int simulation_method; + int binary_mode; + int kernel_id; +} Inputs; + typedef Kokkos::View IntView; typedef Kokkos::View DoubleView; typedef Kokkos::View PointView; @@ -97,7 +110,7 @@ void binary_write( Inputs in, SimulationData SD ); SimulationData binary_read( Inputs in ); // Simulation.c -unsigned long long run_event_based_simulation(Inputs in, SimulationData SD, int mype, double* end, Profile* profile); +unsigned long long run_event_based_simulation(Inputs in, SimulationData SD, int mype, double* end); unsigned long long run_history_based_simulation(Inputs in, SimulationData SD, int mype); KOKKOS_INLINE_FUNCTION void calculate_micro_xs( double p_energy, int nuc, long n_isotopes, diff --git a/kokkos/io.cpp b/kokkos/io.cpp index fdccb4f3..c9f566f1 100644 --- a/kokkos/io.cpp +++ b/kokkos/io.cpp @@ -157,7 +157,6 @@ void print_inputs(Inputs in, int nprocs, int version ) printf("XS Lookups per Particle: "); fancy_int(in.lookups); } printf("Total XS Lookups: "); fancy_int(in.lookups); - printf("Total XS Iterations: "); fancy_int(in.num_iterations); #ifdef MPI printf("MPI Ranks: %d\n", nprocs); printf("Mem Usage per MPI Rank (MB): "); fancy_int(mem_tot); @@ -218,7 +217,6 @@ void print_CLI_error(void) printf(" -h Number of hash bins (only relevant when used with \"-G hash\")\n"); printf(" -b Read or write all data structures to file. If reading, this will skip initialization phase. (read, write)\n"); printf(" -k Specifies which kernel to run. 0 is baseline, 1, 2, etc are optimized variants. (0 is default.)\n"); - printf(" -n Specifies how many kernel iterations to run. (1 is default.)\n"); printf("Default is equivalent to: -m history -s large -l 34 -p 500000 -G unionized -n 1\n"); printf("See readme for full description of default run values\n"); exit(4); @@ -258,9 +256,6 @@ Inputs read_CLI( int argc, char * argv[] ) // defaults to baseline kernel input.kernel_id = 0; - // defaults to one kernel iteration - input.num_iterations = 1; - // defaults to H-M Large benchmark input.HM = (char *) malloc( 6 * sizeof(char) ); input.HM[0] = 'l' ; @@ -276,12 +271,6 @@ Inputs read_CLI( int argc, char * argv[] ) int default_lookups = 1; int default_particles = 1; - // default to zero warmup iterations - input.num_warmups = 1; - - // default to stdout - input.filename = NULL; - // Collect Raw Input for( int i = 1; i < argc; i++ ) { @@ -404,36 +393,6 @@ Inputs read_CLI( int argc, char * argv[] ) else print_CLI_error(); } - // number of kernel iterations (-n) - else if( strcmp(arg, "-n") == 0 ) - { - if( ++i < argc ) - { - input.num_iterations = atoi(argv[i]); - } - else - print_CLI_error(); - } - else if( strcmp(arg, "--csv") == 0 ) - { - if( ++i < argc ) { - input.filename = (char *)malloc(strlen(argv[i]) + 1); - strcpy(input.filename, argv[i]); - } - else - print_CLI_error(); - } - else if( strcmp(arg, "-w") == 0 ) - { - if( ++i < argc) - { - input.num_warmups = atoi(argv[i]); - } - else - print_CLI_error(); - } - else - print_CLI_error(); } // Validate Input @@ -458,10 +417,6 @@ Inputs read_CLI( int argc, char * argv[] ) if( input.hash_bins < 1 ) print_CLI_error(); - // Validate number of iterations - if( input.num_iterations < 1 ) - print_CLI_error(); - // Validate HM size if( strcasecmp(input.HM, "small") != 0 && strcasecmp(input.HM, "large") != 0 && diff --git a/openacc/Main.c b/openacc/Main.c index 5496bb62..2ba76f23 100644 --- a/openacc/Main.c +++ b/openacc/Main.c @@ -68,8 +68,6 @@ int main( int argc, char* argv[] ) border_print(); } - Profile profile; - // Start Simulation Timer omp_start = get_time(); @@ -77,7 +75,7 @@ int main( int argc, char* argv[] ) if( in.simulation_method == EVENT_BASED ) { if( in.kernel_id == 0 ) - verification = run_event_based_simulation(in, SD, mype, &profile); + verification = run_event_based_simulation(in, SD, mype); else { printf("Error: No kernel ID %d found!\n", in.kernel_id); @@ -113,7 +111,5 @@ int main( int argc, char* argv[] ) MPI_Finalize(); #endif - print_profile(profile, in); - return is_invalid_result; } diff --git a/openacc/Simulation.c b/openacc/Simulation.c index 2caa0737..9cb68d8a 100644 --- a/openacc/Simulation.c +++ b/openacc/Simulation.c @@ -11,7 +11,7 @@ // are not yet implemented for this OpenMP targeting offload port. //////////////////////////////////////////////////////////////////////////////////// -unsigned long long run_event_based_simulation(Inputs in, SimulationData SD, int mype, Profile* profile) +unsigned long long run_event_based_simulation(Inputs in, SimulationData SD, int mype) { if( mype == 0) printf("Beginning event based simulation...\n"); @@ -41,8 +41,6 @@ unsigned long long run_event_based_simulation(Inputs in, SimulationData SD, int //////////////////////////////////////////////////////////////////////////////// unsigned long long * verification = (unsigned long long *) malloc(in.lookups * sizeof(unsigned long long)); - double start = get_time(); - int *num_nucs = SD.num_nucs; double *concs = SD.concs; int *mats = SD.mats; @@ -65,80 +63,69 @@ unsigned long long run_event_based_simulation(Inputs in, SimulationData SD, int copyin(index_grid[:SD.length_index_grid]) } - profile->host_to_device_time = get_time() - start; - - int nwarmups = in.num_warmups; - for (int it = 0; it < in.num_iterations + nwarmups; it++) { - if (it == nwarmups) start = get_time(); - #pragma acc parallel loop \ - present(num_nucs[:SD.length_num_nucs], concs[:SD.length_concs], \ - mats[:SD.length_mats], unionized_energy_array[:SD.length_unionized_energy_array], \ - index_grid[:SD.length_index_grid], nuclide_grid[:SD.length_nuclide_grid], verification[:in.lookups]) - for( int i = 0; i < in.lookups; i++ ) + #pragma acc parallel loop \ + present(num_nucs[:SD.length_num_nucs], concs[:SD.length_concs], \ + mats[:SD.length_mats], unionized_energy_array[:SD.length_unionized_energy_array], \ + index_grid[:SD.length_index_grid], nuclide_grid[:SD.length_nuclide_grid], verification[:in.lookups]) + for( int i = 0; i < in.lookups; i++ ) + { + // Set the initial seed value + uint64_t seed = STARTING_SEED; + + // Forward seed to lookup index (we need 2 samples per lookup) + seed = fast_forward_LCG(seed, 2*i); + + // Randomly pick an energy and material for the particle + double p_energy = LCG_random_double(&seed); + int mat = pick_mat(&seed); + + // debugging + //printf("E = %lf mat = %d\n", p_energy, mat); + + double macro_xs_vector[5] = {0}; + + // Perform macroscopic Cross Section Lookup + calculate_macro_xs( + p_energy, // Sampled neutron energy (in lethargy) + mat, // Sampled material type index neutron is in + in.n_isotopes, // Total number of isotopes in simulation + in.n_gridpoints, // Number of gridpoints per isotope in simulation + num_nucs, // 1-D array with number of nuclides per material + concs, // Flattened 2-D array with concentration of each nuclide in each material + unionized_energy_array, // 1-D Unionized energy array + index_grid, // Flattened 2-D grid holding indices into nuclide grid for each unionized energy level + nuclide_grid, // Flattened 2-D grid holding energy levels and XS_data for all nuclides in simulation + mats, // Flattened 2-D array with nuclide indices defining composition of each type of material + macro_xs_vector, // 1-D array with result of the macroscopic cross section (5 different reaction channels) + in.grid_type, // Lookup type (nuclide, hash, or unionized) + in.hash_bins, // Number of hash bins used (if using hash lookup type) + SD.max_num_nucs // Maximum number of nuclides present in any material + ); + + // For verification, and to prevent the compiler from optimizing + // all work out, we interrogate the returned macro_xs_vector array + // to find its maximum value index, then increment the verification + // value by that index. In this implementation, we prevent thread + // contention by using an OMP reduction on the verification value. + // For accelerators, a different approach might be required + // (e.g., atomics, reduction of thread-specific values in large + // array via CUDA thrust, etc). + double max = -1.0; + int max_idx = 0; + for(int j = 0; j < 5; j++ ) { - // Set the initial seed value - uint64_t seed = STARTING_SEED; - - // Forward seed to lookup index (we need 2 samples per lookup) - seed = fast_forward_LCG(seed, 2*i); - - // Randomly pick an energy and material for the particle - double p_energy = LCG_random_double(&seed); - int mat = pick_mat(&seed); - - // debugging - //printf("E = %lf mat = %d\n", p_energy, mat); - - double macro_xs_vector[5] = {0}; - - // Perform macroscopic Cross Section Lookup - calculate_macro_xs( - p_energy, // Sampled neutron energy (in lethargy) - mat, // Sampled material type index neutron is in - in.n_isotopes, // Total number of isotopes in simulation - in.n_gridpoints, // Number of gridpoints per isotope in simulation - num_nucs, // 1-D array with number of nuclides per material - concs, // Flattened 2-D array with concentration of each nuclide in each material - unionized_energy_array, // 1-D Unionized energy array - index_grid, // Flattened 2-D grid holding indices into nuclide grid for each unionized energy level - nuclide_grid, // Flattened 2-D grid holding energy levels and XS_data for all nuclides in simulation - mats, // Flattened 2-D array with nuclide indices defining composition of each type of material - macro_xs_vector, // 1-D array with result of the macroscopic cross section (5 different reaction channels) - in.grid_type, // Lookup type (nuclide, hash, or unionized) - in.hash_bins, // Number of hash bins used (if using hash lookup type) - SD.max_num_nucs // Maximum number of nuclides present in any material - ); - - // For verification, and to prevent the compiler from optimizing - // all work out, we interrogate the returned macro_xs_vector array - // to find its maximum value index, then increment the verification - // value by that index. In this implementation, we prevent thread - // contention by using an OMP reduction on the verification value. - // For accelerators, a different approach might be required - // (e.g., atomics, reduction of thread-specific values in large - // array via CUDA thrust, etc). - double max = -1.0; - int max_idx = 0; - for(int j = 0; j < 5; j++ ) + if( macro_xs_vector[j] > max ) { - if( macro_xs_vector[j] > max ) - { - max = macro_xs_vector[j]; - max_idx = j; - } + max = macro_xs_vector[j]; + max_idx = j; } - verification[i] = max_idx+1; } + verification[i] = max_idx+1; } - profile->kernel_time = get_time() - start; - - start = get_time(); #pragma acc exit data \ copyout(verification[0:in.lookups]) - profile->device_to_host_time = get_time() - start; - // Reduce validation hash on the host unsigned long long validation_hash = 0; for( int i = 0; i < in.lookups; i++ ) diff --git a/openacc/XSbench_header.h b/openacc/XSbench_header.h index b590c333..c19e3ae5 100644 --- a/openacc/XSbench_header.h +++ b/openacc/XSbench_header.h @@ -15,7 +15,6 @@ #include #include #include -#include "../XSbench_shared_header.h" // Papi Header #ifdef PAPI @@ -49,6 +48,20 @@ typedef struct{ double nu_fission_xs; } NuclideGridPoint; +typedef struct{ + int nthreads; + long n_isotopes; + long n_gridpoints; + int lookups; + char * HM; + int grid_type; // 0: Unionized Grid (default) 1: Nuclide Grid + int hash_bins; + int particles; + int simulation_method; + int binary_mode; + int kernel_id; +} Inputs; + typedef struct{ int * num_nucs; // Length = length_num_nucs; double * concs; // Length = length_concs @@ -82,7 +95,7 @@ void binary_write( Inputs in, SimulationData SD ); SimulationData binary_read( Inputs in ); // Simulation.c -unsigned long long run_event_based_simulation(Inputs in, SimulationData SD, int mype, Profile* profile); +unsigned long long run_event_based_simulation(Inputs in, SimulationData SD, int mype); unsigned long long run_history_based_simulation(Inputs in, SimulationData SD, int mype); #pragma acc routine seq void calculate_micro_xs( double p_energy, int nuc, long n_isotopes, diff --git a/openacc/io.c b/openacc/io.c index 12d03050..884e95f2 100644 --- a/openacc/io.c +++ b/openacc/io.c @@ -157,7 +157,6 @@ void print_inputs(Inputs in, int nprocs, int version ) printf("XS Lookups per Particle: "); fancy_int(in.lookups); } printf("Total XS Lookups: "); fancy_int(in.lookups); - printf("Total XS Iterations: "); fancy_int(in.num_iterations); #ifdef MPI printf("MPI Ranks: %d\n", nprocs); printf("Mem Usage per MPI Rank (MB): "); fancy_int(mem_tot); @@ -218,7 +217,6 @@ void print_CLI_error(void) printf(" -h Number of hash bins (only relevant when used with \"-G hash\")\n"); printf(" -b Read or write all data structures to file. If reading, this will skip initialization phase. (read, write)\n"); printf(" -k Specifies which kernel to run. 0 is baseline, 1, 2, etc are optimized variants. (0 is default.)\n"); - printf(" -n Specifies how many kernel iterations to run. (1 is default.)\n"); printf("Default is equivalent to: -m history -s large -l 34 -p 500000 -G unionized -n 1\n"); printf("See readme for full description of default run values\n"); exit(4); @@ -258,15 +256,6 @@ Inputs read_CLI( int argc, char * argv[] ) // defaults to baseline kernel input.kernel_id = 0; - // default to one kernel iteration - input.num_iterations = 1; - // - // default to zero warmup iterations - input.num_warmups = 1; - - // default to stdout - input.filename = NULL; - // defaults to H-M Large benchmark input.HM = (char *) malloc( 6 * sizeof(char) ); input.HM[0] = 'l' ; @@ -404,36 +393,6 @@ Inputs read_CLI( int argc, char * argv[] ) else print_CLI_error(); } - // number of kernel iterations (-n) - else if( strcmp(arg, "-n") == 0 ) - { - if( ++i < argc) - { - input.num_iterations = atoi(argv[i]); - } - else - print_CLI_error(); - } - else if( strcmp(arg, "--csv") == 0 ) - { - if( ++i < argc ) { - input.filename = (char *)malloc(strlen(argv[i]) + 1); - strcpy(input.filename, argv[i]); - } - else - print_CLI_error(); - } - else if( strcmp(arg, "-w") == 0 ) - { - if( ++i < argc) - { - input.num_warmups = atoi(argv[i]); - } - else - print_CLI_error(); - } - else - print_CLI_error(); } // Validate Input @@ -458,10 +417,6 @@ Inputs read_CLI( int argc, char * argv[] ) if( input.hash_bins < 1 ) print_CLI_error(); - // Validate number of iterations - if ( input.num_iterations < 1 ) - print_CLI_error(); - // Validate HM size if( strcasecmp(input.HM, "small") != 0 && strcasecmp(input.HM, "large") != 0 && diff --git a/openmp-offload/Main.c b/openmp-offload/Main.c index cabac5f3..cd631e8f 100644 --- a/openmp-offload/Main.c +++ b/openmp-offload/Main.c @@ -68,8 +68,6 @@ int main( int argc, char* argv[] ) border_print(); } - Profile profile; - // Start Simulation Timer omp_start = omp_get_wtime(); @@ -77,7 +75,7 @@ int main( int argc, char* argv[] ) if( in.simulation_method == EVENT_BASED ) { if( in.kernel_id == 0 ) - verification = run_event_based_simulation(in, SD, mype, &profile); + verification = run_event_based_simulation(in, SD, mype); else { printf("Error: No kernel ID %d found!\n", in.kernel_id); @@ -113,7 +111,5 @@ int main( int argc, char* argv[] ) MPI_Finalize(); #endif - print_profile(profile, in); - return is_invalid_result; } diff --git a/openmp-offload/Simulation.c b/openmp-offload/Simulation.c index 34794cad..12ff0f1a 100644 --- a/openmp-offload/Simulation.c +++ b/openmp-offload/Simulation.c @@ -11,7 +11,7 @@ // are not yet implemented for this OpenMP targeting offload port. //////////////////////////////////////////////////////////////////////////////////// -unsigned long long run_event_based_simulation(Inputs in, SimulationData SD, int mype, Profile* profile) +unsigned long long run_event_based_simulation(Inputs in, SimulationData SD, int mype) { if( mype == 0) printf("Beginning event based simulation...\n"); @@ -41,8 +41,6 @@ unsigned long long run_event_based_simulation(Inputs in, SimulationData SD, int //////////////////////////////////////////////////////////////////////////////// unsigned long long * verification = (unsigned long long *) malloc(in.lookups * sizeof(unsigned long long)); - double start = get_time(); - int *num_nucs = SD.num_nucs; double *concs = SD.concs; int *mats = SD.mats; @@ -63,77 +61,66 @@ unsigned long long run_event_based_simulation(Inputs in, SimulationData SD, int #pragma omp target enter data map(to: index_grid[:SD.length_index_grid]) } - profile->host_to_device_time = get_time() - start; - - int nwarmups = in.num_warmups; - for (int it = 0; it < in.num_iterations + nwarmups; it++) { - if (it == nwarmups) start = get_time(); - #pragma omp target teams distribute parallel for - for( int i = 0; i < in.lookups; i++ ) + #pragma omp target teams distribute parallel for + for( int i = 0; i < in.lookups; i++ ) + { + // Set the initial seed value + uint64_t seed = STARTING_SEED; + + // Forward seed to lookup index (we need 2 samples per lookup) + seed = fast_forward_LCG(seed, 2*i); + + // Randomly pick an energy and material for the particle + double p_energy = LCG_random_double(&seed); + int mat = pick_mat(&seed); + + // debugging + //printf("E = %lf mat = %d\n", p_energy, mat); + + double macro_xs_vector[5] = {0}; + + // Perform macroscopic Cross Section Lookup + calculate_macro_xs( + p_energy, // Sampled neutron energy (in lethargy) + mat, // Sampled material type index neutron is in + in.n_isotopes, // Total number of isotopes in simulation + in.n_gridpoints, // Number of gridpoints per isotope in simulation + num_nucs, // 1-D array with number of nuclides per material + concs, // Flattened 2-D array with concentration of each nuclide in each material + unionized_energy_array, // 1-D Unionized energy array + index_grid, // Flattened 2-D grid holding indices into nuclide grid for each unionized energy level + nuclide_grid, // Flattened 2-D grid holding energy levels and XS_data for all nuclides in simulation + mats, // Flattened 2-D array with nuclide indices defining composition of each type of material + macro_xs_vector, // 1-D array with result of the macroscopic cross section (5 different reaction channels) + in.grid_type, // Lookup type (nuclide, hash, or unionized) + in.hash_bins, // Number of hash bins used (if using hash lookup type) + SD.max_num_nucs // Maximum number of nuclides present in any material + ); + + // For verification, and to prevent the compiler from optimizing + // all work out, we interrogate the returned macro_xs_vector array + // to find its maximum value index, then increment the verification + // value by that index. In this implementation, we prevent thread + // contention by using an OMP reduction on the verification value. + // For accelerators, a different approach might be required + // (e.g., atomics, reduction of thread-specific values in large + // array via CUDA thrust, etc). + double max = -1.0; + int max_idx = 0; + for(int j = 0; j < 5; j++ ) { - // Set the initial seed value - uint64_t seed = STARTING_SEED; - - // Forward seed to lookup index (we need 2 samples per lookup) - seed = fast_forward_LCG(seed, 2*i); - - // Randomly pick an energy and material for the particle - double p_energy = LCG_random_double(&seed); - int mat = pick_mat(&seed); - - // debugging - //printf("E = %lf mat = %d\n", p_energy, mat); - - double macro_xs_vector[5] = {0}; - - // Perform macroscopic Cross Section Lookup - calculate_macro_xs( - p_energy, // Sampled neutron energy (in lethargy) - mat, // Sampled material type index neutron is in - in.n_isotopes, // Total number of isotopes in simulation - in.n_gridpoints, // Number of gridpoints per isotope in simulation - num_nucs, // 1-D array with number of nuclides per material - concs, // Flattened 2-D array with concentration of each nuclide in each material - unionized_energy_array, // 1-D Unionized energy array - index_grid, // Flattened 2-D grid holding indices into nuclide grid for each unionized energy level - nuclide_grid, // Flattened 2-D grid holding energy levels and XS_data for all nuclides in simulation - mats, // Flattened 2-D array with nuclide indices defining composition of each type of material - macro_xs_vector, // 1-D array with result of the macroscopic cross section (5 different reaction channels) - in.grid_type, // Lookup type (nuclide, hash, or unionized) - in.hash_bins, // Number of hash bins used (if using hash lookup type) - SD.max_num_nucs // Maximum number of nuclides present in any material - ); - - // For verification, and to prevent the compiler from optimizing - // all work out, we interrogate the returned macro_xs_vector array - // to find its maximum value index, then increment the verification - // value by that index. In this implementation, we prevent thread - // contention by using an OMP reduction on the verification value. - // For accelerators, a different approach might be required - // (e.g., atomics, reduction of thread-specific values in large - // array via CUDA thrust, etc). - double max = -1.0; - int max_idx = 0; - for(int j = 0; j < 5; j++ ) + if( macro_xs_vector[j] > max ) { - if( macro_xs_vector[j] > max ) - { - max = macro_xs_vector[j]; - max_idx = j; - } + max = macro_xs_vector[j]; + max_idx = j; } - verification[i] = max_idx+1; } + verification[i] = max_idx+1; } - profile->kernel_time = get_time() - start; - - start = get_time(); #pragma omp target exit data\ map(from: verification[:in.lookups]) - profile->device_to_host_time = get_time() - start; - // Reduce validation hash on the host unsigned long long validation_hash = 0; for( int i = 0; i < in.lookups; i++ ) diff --git a/openmp-offload/XSbench_header.h b/openmp-offload/XSbench_header.h index c31237f1..8a4ea3ce 100644 --- a/openmp-offload/XSbench_header.h +++ b/openmp-offload/XSbench_header.h @@ -12,7 +12,6 @@ #include #include #include -#include "../XSbench_shared_header.h" // Papi Header #ifdef PAPI @@ -46,6 +45,20 @@ typedef struct{ double nu_fission_xs; } NuclideGridPoint; +typedef struct{ + int nthreads; + long n_isotopes; + long n_gridpoints; + int lookups; + char * HM; + int grid_type; // 0: Unionized Grid (default) 1: Nuclide Grid + int hash_bins; + int particles; + int simulation_method; + int binary_mode; + int kernel_id; +} Inputs; + typedef struct{ int * num_nucs; // Length = length_num_nucs; double * concs; // Length = length_concs @@ -79,7 +92,7 @@ void binary_write( Inputs in, SimulationData SD ); SimulationData binary_read( Inputs in ); // Simulation.c -unsigned long long run_event_based_simulation(Inputs in, SimulationData SD, int mype, Profile* profile); +unsigned long long run_event_based_simulation(Inputs in, SimulationData SD, int mype); unsigned long long run_history_based_simulation(Inputs in, SimulationData SD, int mype); void calculate_micro_xs( double p_energy, int nuc, long n_isotopes, long n_gridpoints, diff --git a/openmp-offload/io.c b/openmp-offload/io.c index 0236c151..e4f3c1d5 100644 --- a/openmp-offload/io.c +++ b/openmp-offload/io.c @@ -157,7 +157,6 @@ void print_inputs(Inputs in, int nprocs, int version ) printf("XS Lookups per Particle: "); fancy_int(in.lookups); } printf("Total XS Lookups: "); fancy_int(in.lookups); - printf("Total XS Iterations: "); fancy_int(in.num_iterations); #ifdef MPI printf("MPI Ranks: %d\n", nprocs); printf("Mem Usage per MPI Rank (MB): "); fancy_int(mem_tot); @@ -218,7 +217,6 @@ void print_CLI_error(void) printf(" -h Number of hash bins (only relevant when used with \"-G hash\")\n"); printf(" -b Read or write all data structures to file. If reading, this will skip initialization phase. (read, write)\n"); printf(" -k Specifies which kernel to run. 0 is baseline, 1, 2, etc are optimized variants. (0 is default.)\n"); - printf(" -n Specifies how many kernel iterations to run. (1 is default.)\n"); printf("Default is equivalent to: -m history -s large -l 34 -p 500000 -G unionized -n 1\n"); printf("See readme for full description of default run values\n"); exit(4); @@ -258,15 +256,6 @@ Inputs read_CLI( int argc, char * argv[] ) // defaults to baseline kernel input.kernel_id = 0; - // default to one kernel iteration - input.num_iterations = 1; - - // default to zero warmup iterations - input.num_warmups = 1; - - // default to stdout - input.filename = NULL; - // defaults to H-M Large benchmark input.HM = (char *) malloc( 6 * sizeof(char) ); input.HM[0] = 'l' ; @@ -404,36 +393,6 @@ Inputs read_CLI( int argc, char * argv[] ) else print_CLI_error(); } - // number of kernel iterations (-n) - else if( strcmp(arg, "-n") == 0 ) - { - if( ++i < argc) - { - input.num_iterations = atoi(argv[i]); - } - else - print_CLI_error(); - } - else if( strcmp(arg, "--csv") == 0 ) - { - if( ++i < argc ) { - input.filename = (char *)malloc(strlen(argv[i]) + 1); - strcpy(input.filename, argv[i]); - } - else - print_CLI_error(); - } - else if( strcmp(arg, "-w") == 0 ) - { - if( ++i < argc) - { - input.num_warmups = atoi(argv[i]); - } - else - print_CLI_error(); - } - else - print_CLI_error(); } // Validate Input @@ -458,10 +417,6 @@ Inputs read_CLI( int argc, char * argv[] ) if( input.hash_bins < 1 ) print_CLI_error(); - // Validate number of iterations - if ( input.num_iterations < 1 ) - print_CLI_error(); - // Validate HM size if( strcasecmp(input.HM, "small") != 0 && strcasecmp(input.HM, "large") != 0 && diff --git a/openmp-threading/XSbench_header.h b/openmp-threading/XSbench_header.h index 53157a66..eccbac68 100644 --- a/openmp-threading/XSbench_header.h +++ b/openmp-threading/XSbench_header.h @@ -60,6 +60,20 @@ typedef struct{ double nu_fission_xs; } NuclideGridPoint; +typedef struct{ + int nthreads; + long n_isotopes; + long n_gridpoints; + int lookups; + char * HM; + int grid_type; // 0: Unionized Grid (default) 1: Nuclide Grid + int hash_bins; + int particles; + int simulation_method; + int binary_mode; + int kernel_id; +} Inputs; + typedef struct{ int nthreads; long n_isotopes; diff --git a/raja/CMakeLists.txt b/raja/CMakeLists.txt index e744b25c..1ba9f828 100644 --- a/raja/CMakeLists.txt +++ b/raja/CMakeLists.txt @@ -25,14 +25,16 @@ find_package(RAJA REQUIRED) find_package(Umpire REQUIRED) target_link_libraries(XSBench RAJA umpire) -if (${RAJA_BACK_END} STREQUAL "CUDA") +if (RAJA_ENABLE_CUDA) + message(STATUS "RAJA CUDA support enabled") enable_language(CUDA) set(CMAKE_CUDA_STANDARD 14) set_source_files_properties(${SOURCE} PROPERTIES LANGUAGE CUDA) set(CMAKE_CUDA_FLAGS "${CMAKE_CUDA_FLAGS} -extended-lambda --expt-relaxed-constexpr --restrict --keep") set_target_properties(XSBench PROPERTIES CUDA_RESOLVE_DEVICE_SYMBOLS ON) -elseif (${RAJA_BACK_END} STREQUAL "HIP") + elseif (RAJA_ENABLE_HIP) + message(STATUS "RAJA HIP support enabled") enable_language(HIP) set(CMAKE_HIP_STANDARD 14) diff --git a/raja/Main.cpp b/raja/Main.cpp index 9cf7ca0d..16083609 100644 --- a/raja/Main.cpp +++ b/raja/Main.cpp @@ -55,13 +55,11 @@ int main( int argc, char* argv[] ) // Start Simulation Timer omp_start = get_time(); - Profile profile; - // Run simulation if( in.simulation_method == EVENT_BASED ) { if( in.kernel_id == 0 ) - verification = run_event_based_simulation_baseline(in, SD, mype, &profile); + verification = run_event_based_simulation_baseline(in, SD, mype); else { printf("Error: No kernel ID %d found!\n", in.kernel_id); exit(1); @@ -90,7 +88,5 @@ int main( int argc, char* argv[] ) // Print / Save Results and Exit int is_invalid_result = print_results( in, mype, omp_end-omp_start, nprocs, verification ); - print_profile(profile, in); - return is_invalid_result; } diff --git a/raja/Simulation.cpp b/raja/Simulation.cpp index ed52bb03..d19985c2 100644 --- a/raja/Simulation.cpp +++ b/raja/Simulation.cpp @@ -22,7 +22,7 @@ using policy = RAJA::omp_parallel_for_exec; using policy = RAJA::seq_exec; #endif -unsigned long long run_event_based_simulation_baseline(Inputs in, SimulationData SD, int mype, Profile* profile) +unsigned long long run_event_based_simulation_baseline(Inputs in, SimulationData SD, int mype) { //////////////////////////////////////////////////////////////////////////////// // Configure & Launch Simulation Kernel @@ -32,74 +32,63 @@ unsigned long long run_event_based_simulation_baseline(Inputs in, SimulationData auto& rm = umpire::ResourceManager::getInstance(); - double start = get_time(); SimulationData GSD = move_simulation_data_to_device(in, mype, SD); - profile->host_to_device_time = get_time() - start; int nthreads = 256; int nblocks = ceil( (double) in.lookups / (double) nthreads); - int nwarmups = in.num_warmups; start = 0.0; - for (int it = 0; it < in.num_iterations + nwarmups; it++) { - if (it == nwarmups) { - start = get_time(); - } - RAJA::forall(RAJA::RangeSegment(0, in.lookups), [=] RAJA_HOST_DEVICE (int i) { - // Set the initial seed value - uint64_t seed = STARTING_SEED; - - // Forward seed to lookup index (we need 2 samples per lookup) - seed = fast_forward_LCG(seed, 2*i); - - // Randomly pick an energy and material for the particle - double p_energy = LCG_random_double(&seed); - int mat = pick_mat(&seed); - - double macro_xs_vector[5] = {0}; - - // Perform macroscopic Cross Section Lookup - calculate_macro_xs( - p_energy, // Sampled neutron energy (in lethargy) - mat, // Sampled material type index neutron is in - in.n_isotopes, // Total number of isotopes in simulation - in.n_gridpoints, // Number of gridpoints per isotope in simulation - GSD.num_nucs, // 1-D array with number of nuclides per material - GSD.concs, // Flattened 2-D array with concentration of each nuclide in each material - GSD.unionized_energy_array, // 1-D Unionized energy array - GSD.index_grid, // Flattened 2-D grid holding indices into nuclide grid for each unionized energy level - GSD.nuclide_grid, // Flattened 2-D grid holding energy levels and XS_data for all nuclides in simulation - GSD.mats, // Flattened 2-D array with nuclide indices defining composition of each type of material - macro_xs_vector, // 1-D array with result of the macroscopic cross section (5 different reaction channels) - in.grid_type, // Lookup type (nuclide, hash, or unionized) - in.hash_bins, // Number of hash bins used (if using hash lookup type) - GSD.max_num_nucs // Maximum number of nuclides present in any material - ); - - - double max = -1.0; - int max_idx = 0; - for(int j = 0; j < 5; j++ ) - { - if( macro_xs_vector[j] > max ) - { - max = macro_xs_vector[j]; - max_idx = j; - } - } - GSD.verification[i] = max_idx+1; - }); - } - profile->kernel_time = get_time() - start; + RAJA::forall(RAJA::RangeSegment(0, in.lookups), [=] RAJA_HOST_DEVICE (int i) { + // Set the initial seed value + uint64_t seed = STARTING_SEED; + + // Forward seed to lookup index (we need 2 samples per lookup) + seed = fast_forward_LCG(seed, 2*i); + + // Randomly pick an energy and material for the particle + double p_energy = LCG_random_double(&seed); + int mat = pick_mat(&seed); + + double macro_xs_vector[5] = {0}; + + // Perform macroscopic Cross Section Lookup + calculate_macro_xs( + p_energy, // Sampled neutron energy (in lethargy) + mat, // Sampled material type index neutron is in + in.n_isotopes, // Total number of isotopes in simulation + in.n_gridpoints, // Number of gridpoints per isotope in simulation + GSD.num_nucs, // 1-D array with number of nuclides per material + GSD.concs, // Flattened 2-D array with concentration of each nuclide in each material + GSD.unionized_energy_array, // 1-D Unionized energy array + GSD.index_grid, // Flattened 2-D grid holding indices into nuclide grid for each unionized energy level + GSD.nuclide_grid, // Flattened 2-D grid holding energy levels and XS_data for all nuclides in simulation + GSD.mats, // Flattened 2-D array with nuclide indices defining composition of each type of material + macro_xs_vector, // 1-D array with result of the macroscopic cross section (5 different reaction channels) + in.grid_type, // Lookup type (nuclide, hash, or unionized) + in.hash_bins, // Number of hash bins used (if using hash lookup type) + GSD.max_num_nucs // Maximum number of nuclides present in any material + ); + + + double max = -1.0; + int max_idx = 0; + for(int j = 0; j < 5; j++ ) + { + if( macro_xs_vector[j] > max ) + { + max = macro_xs_vector[j]; + max_idx = j; + } + } + GSD.verification[i] = max_idx+1; + }); //////////////////////////////////////////////////////////////////////////////// // Reduce Verification Results //////////////////////////////////////////////////////////////////////////////// if( mype == 0) printf("Reducing verification results...\n"); - start = get_time(); rm.copy(SD.verification, GSD.verification); - profile->device_to_host_time = get_time() - start; unsigned long long verification_scalar = 0; for(int i = 0; i < in.lookups; i++ ) diff --git a/raja/XSbench_header.h b/raja/XSbench_header.h index d57a4b33..3146b3f2 100644 --- a/raja/XSbench_header.h +++ b/raja/XSbench_header.h @@ -7,7 +7,6 @@ #include #include #include -#include "../XSbench_shared_header.h" #include #include "umpire/Allocator.hpp" @@ -63,6 +62,20 @@ typedef struct{ double nu_fission_xs; } NuclideGridPoint; +typedef struct{ + int nthreads; + long n_isotopes; + long n_gridpoints; + int lookups; + char * HM; + int grid_type; // 0: Unionized Grid (default) 1: Nuclide Grid + int hash_bins; + int particles; + int simulation_method; + int binary_mode; + int kernel_id; +} Inputs; + typedef struct{ int * num_nucs; // Length = length_num_nucs; double * concs; // Length = length_concs @@ -98,7 +111,7 @@ void binary_write( Inputs in, SimulationData SD ); SimulationData binary_read( Inputs in ); // Simulation.cpp -unsigned long long run_event_based_simulation_baseline(Inputs in, SimulationData SD, int mype, Profile* profile); +unsigned long long run_event_based_simulation_baseline(Inputs in, SimulationData SD, int mype); RAJA_HOST_DEVICE void calculate_micro_xs( double p_energy, int nuc, long n_isotopes, long n_gridpoints, double * __restrict__ egrid, int * __restrict__ index_data, diff --git a/raja/io.cpp b/raja/io.cpp index ae93c45e..759915a6 100644 --- a/raja/io.cpp +++ b/raja/io.cpp @@ -76,7 +76,6 @@ int print_results( Inputs in, int mype, double runtime, int nprocs, printf("Lookups/s: "); fancy_int(lookups_per_sec); #endif - printf("Total XS Iterations: "); fancy_int(in.num_iterations); } unsigned long long large = 0; @@ -228,7 +227,6 @@ void print_CLI_error(void) printf(" -h Number of hash bins (only relevant when used with \"-G hash\")\n"); printf(" -b Read or write all data structures to file. If reading, this will skip initialization phase. (read, write)\n"); printf(" -k Specifies which kernel to run. 0 is baseline, 1, 2, etc are optimized variants. (0 is default.)\n"); - printf(" -n Specifies how many kernel iterations to run. (1 is default.)\n"); printf("Default is equivalent to: -m history -s large -l 34 -p 500000 -G unionized -k 0 -n 1\n"); printf("See readme for full description of default run values\n"); exit(4); @@ -272,15 +270,6 @@ Inputs read_CLI( int argc, char * argv[] ) // defaults to baseline kernel input.kernel_id = 0; - // defaults to one kernel iteration - input.num_iterations = 1; - - // default to zero warmup iterations - input.num_warmups = 1; - - // default to stdout - input.filename = NULL; - // defaults to H-M Large benchmark input.HM = (char *) malloc( 6 * sizeof(char) ); input.HM[0] = 'l' ; @@ -418,36 +407,6 @@ Inputs read_CLI( int argc, char * argv[] ) else print_CLI_error(); } - // number of kernel iterations (-n) - else if( strcmp(arg, "-n") == 0 ) - { - if( ++i < argc ) - { - input.num_iterations = atoi(argv[i]); - } - else - print_CLI_error(); - } - else if( strcmp(arg, "--csv") == 0 ) - { - if( ++i < argc ) { - input.filename = (char *)malloc(strlen(argv[i]) + 1); - strcpy(input.filename, argv[i]); - } - else - print_CLI_error(); - } - else if( strcmp(arg, "-w") == 0 ) - { - if( ++i < argc) - { - input.num_warmups = atoi(argv[i]); - } - else - print_CLI_error(); - } - else - print_CLI_error(); } // Validate Input @@ -472,10 +431,6 @@ Inputs read_CLI( int argc, char * argv[] ) if( input.hash_bins < 1 ) print_CLI_error(); - // Validate number of iterations - if( input.num_iterations < 1 ) - print_CLI_error(); - // Validate HM size if( strcasecmp(input.HM, "small") != 0 && strcasecmp(input.HM, "large") != 0 && diff --git a/sycl-acc/Main.cpp b/sycl-acc/Main.cpp index 22d44832..0c475947 100644 --- a/sycl-acc/Main.cpp +++ b/sycl-acc/Main.cpp @@ -68,8 +68,6 @@ int main( int argc, char* argv[] ) border_print(); } - Profile profile; - // Start Simulation Timer omp_start = get_time(); double kernel_init_time; @@ -79,7 +77,7 @@ int main( int argc, char* argv[] ) { if( in.kernel_id == 0 ) { - verification = run_event_based_simulation(in, SD, mype, &kernel_init_time, &profile); + verification = run_event_based_simulation(in, SD, mype, &kernel_init_time); } else { @@ -117,7 +115,5 @@ int main( int argc, char* argv[] ) MPI_Finalize(); #endif - print_profile(profile, in); - return is_invalid_result; } diff --git a/sycl-acc/Simulation.cpp b/sycl-acc/Simulation.cpp index f85cc5a7..ad56887d 100644 --- a/sycl-acc/Simulation.cpp +++ b/sycl-acc/Simulation.cpp @@ -12,7 +12,7 @@ //////////////////////////////////////////////////////////////////////////////////// // use SYCL namespace to reduce symbol names -unsigned long long run_event_based_simulation(Inputs in, SimulationData SD, int mype, double * kernel_init_time, Profile* profile) { +unsigned long long run_event_based_simulation(Inputs in, SimulationData SD, int mype, double * kernel_init_time) { //////////////////////////////////////////////////////////////////////////////// // SUMMARY: Simulation Data Structure Manifest for "SD" Object @@ -51,7 +51,6 @@ unsigned long long run_event_based_simulation(Inputs in, SimulationData SD, int // Create Device Buffers //////////////////////////////////////////////////////////////////////////////// - double startP = get_time(); // assign SYCL buffer to existing memory sycl::buffer num_nucs_d(SD.num_nucs,SD.length_num_nucs); sycl::buffer concs_d(SD.concs, SD.length_concs); @@ -60,105 +59,92 @@ unsigned long long run_event_based_simulation(Inputs in, SimulationData SD, int sycl::buffer index_grid_d(SD.index_grid, SD.length_index_grid); sycl::buffer nuclide_grid_d(SD.nuclide_grid, SD.length_nuclide_grid); sycl::buffer verification_d(verification_host, in.lookups); - profile->host_to_device_time = get_time() - startP; if(mype==0) printf("Beginning event based simulation...\n"); - int nwarmups = in.num_warmups; - startP = 0.0; - for (int it = 0; it < in.num_iterations + nwarmups; it++) { - if (it == nwarmups) { - sycl_q.wait(); - startP = get_time(); - } + //////////////////////////////////////////////////////////////////////////////// + // Define Device Kernel + //////////////////////////////////////////////////////////////////////////////// + // queue a kernel to be run, as a lambda + sycl_q.submit([&](sycl::handler &cgh) { //////////////////////////////////////////////////////////////////////////////// - // Define Device Kernel + // Create Device Accessors for Device Buffers //////////////////////////////////////////////////////////////////////////////// + sycl::accessor num_nucs {num_nucs_d, cgh, sycl::read_only}; + sycl::accessor concs {concs_d, cgh, sycl::read_only}; + sycl::accessor mats {mats_d, cgh, sycl::read_only}; + sycl::accessor unionized_energy_array {unionized_energy_array_d, cgh, sycl::read_only}; + sycl::accessor index_grid {index_grid_d, cgh, sycl::read_only}; + sycl::accessor nuclide_grid {nuclide_grid_d, cgh, sycl::read_only}; + sycl::accessor verification {verification_d, cgh, sycl::write_only, sycl::no_init}; - // queue a kernel to be run, as a lambda - sycl_q.submit([&](sycl::handler &cgh) { - //////////////////////////////////////////////////////////////////////////////// - // Create Device Accessors for Device Buffers - //////////////////////////////////////////////////////////////////////////////// - sycl::accessor num_nucs {num_nucs_d, cgh, sycl::read_only}; - sycl::accessor concs {concs_d, cgh, sycl::read_only}; - sycl::accessor mats {mats_d, cgh, sycl::read_only}; - sycl::accessor unionized_energy_array {unionized_energy_array_d, cgh, sycl::read_only}; - sycl::accessor index_grid {index_grid_d, cgh, sycl::read_only}; - sycl::accessor nuclide_grid {nuclide_grid_d, cgh, sycl::read_only}; - sycl::accessor verification {verification_d, cgh, sycl::write_only, sycl::no_init}; - - //////////////////////////////////////////////////////////////////////////////// - // XS Lookup Simulation Loop - //////////////////////////////////////////////////////////////////////////////// - cgh.parallel_for(sycl::range<1>(in.lookups), [=](sycl::id<1> idx) { - // get the index to operate on, first dimemsion - size_t i = idx[0]; - - // Set the initial seed value - uint64_t seed = STARTING_SEED; - - // Forward seed to lookup index (we need 2 samples per lookup) - seed = fast_forward_LCG(seed, 2*i); - - // Randomly pick an energy and material for the particle - double p_energy = LCG_random_double(&seed); - int mat = pick_mat(&seed); - - // debugging - //printf("E = %lf mat = %d\n", p_energy, mat); - - double macro_xs_vector[5] = {0}; - - // Perform macroscopic Cross Section Lookup - calculate_macro_xs( - p_energy, // Sampled neutron energy (in lethargy) - mat, // Sampled material type index neutron is in - in.n_isotopes, // Total number of isotopes in simulation - in.n_gridpoints, // Number of gridpoints per isotope in simulation - num_nucs, // 1-D array with number of nuclides per material - concs, // Flattened 2-D array with concentration of each nuclide in each material - unionized_energy_array, // 1-D Unionized energy array - index_grid, // Flattened 2-D grid holding indices into nuclide grid for each unionized energy level - nuclide_grid, // Flattened 2-D grid holding energy levels and XS_data for all nuclides in simulation - mats, // Flattened 2-D array with nuclide indices defining composition of each type of material - macro_xs_vector, // 1-D array with result of the macroscopic cross section (5 different reaction channels) - in.grid_type, // Lookup type (nuclide, hash, or unionized) - in.hash_bins, // Number of hash bins used (if using hash lookup type) - SD.max_num_nucs // Maximum number of nuclides present in any material - ); - - // For verification, and to prevent the compiler from optimizing - // all work out, we interrogate the returned macro_xs_vector array - // to find its maximum value index, then increment the verification - // value by that index. In this implementation, we store to a global - // array that will get tranferred back and reduced on the host. - double max = -1.0; - int max_idx = 0; - for(int j = 0; j < 5; j++ ) + //////////////////////////////////////////////////////////////////////////////// + // XS Lookup Simulation Loop + //////////////////////////////////////////////////////////////////////////////// + cgh.parallel_for(sycl::range<1>(in.lookups), [=](sycl::id<1> idx) { + // get the index to operate on, first dimemsion + size_t i = idx[0]; + + // Set the initial seed value + uint64_t seed = STARTING_SEED; + + // Forward seed to lookup index (we need 2 samples per lookup) + seed = fast_forward_LCG(seed, 2*i); + + // Randomly pick an energy and material for the particle + double p_energy = LCG_random_double(&seed); + int mat = pick_mat(&seed); + + // debugging + //printf("E = %lf mat = %d\n", p_energy, mat); + + double macro_xs_vector[5] = {0}; + + // Perform macroscopic Cross Section Lookup + calculate_macro_xs( + p_energy, // Sampled neutron energy (in lethargy) + mat, // Sampled material type index neutron is in + in.n_isotopes, // Total number of isotopes in simulation + in.n_gridpoints, // Number of gridpoints per isotope in simulation + num_nucs, // 1-D array with number of nuclides per material + concs, // Flattened 2-D array with concentration of each nuclide in each material + unionized_energy_array, // 1-D Unionized energy array + index_grid, // Flattened 2-D grid holding indices into nuclide grid for each unionized energy level + nuclide_grid, // Flattened 2-D grid holding energy levels and XS_data for all nuclides in simulation + mats, // Flattened 2-D array with nuclide indices defining composition of each type of material + macro_xs_vector, // 1-D array with result of the macroscopic cross section (5 different reaction channels) + in.grid_type, // Lookup type (nuclide, hash, or unionized) + in.hash_bins, // Number of hash bins used (if using hash lookup type) + SD.max_num_nucs // Maximum number of nuclides present in any material + ); + + // For verification, and to prevent the compiler from optimizing + // all work out, we interrogate the returned macro_xs_vector array + // to find its maximum value index, then increment the verification + // value by that index. In this implementation, we store to a global + // array that will get tranferred back and reduced on the host. + double max = -1.0; + int max_idx = 0; + for(int j = 0; j < 5; j++ ) + { + if( macro_xs_vector[j] > max ) { - if( macro_xs_vector[j] > max ) - { - max = macro_xs_vector[j]; - max_idx = j; - } + max = macro_xs_vector[j]; + max_idx = j; } - verification[i] = max_idx+1; + } + verification[i] = max_idx+1; - }); }); - } + }); stop = get_time(); sycl_q.wait(); - profile->kernel_time = get_time() - startP; if(mype==0) printf("Kernel initialization, compilation, and launch took %.2lf seconds.\n", stop-start); - startP = get_time(); verification_d.get_host_access(); - profile->device_to_host_time = get_time() - startP; // Host reduces the verification array unsigned long long verification_scalar = 0; diff --git a/sycl-acc/XSbench_header.h b/sycl-acc/XSbench_header.h index 5e492ee1..b4ad5d1b 100644 --- a/sycl-acc/XSbench_header.h +++ b/sycl-acc/XSbench_header.h @@ -11,7 +11,7 @@ #include #include #include -#include "../XSbench_shared_header.h" + // Papi Header #ifdef PAPI @@ -45,6 +45,20 @@ typedef struct{ double nu_fission_xs; } NuclideGridPoint; +typedef struct{ + int nthreads; + long n_isotopes; + long n_gridpoints; + int lookups; + char * HM; + int grid_type; // 0: Unionized Grid (default) 1: Nuclide Grid + int hash_bins; + int particles; + int simulation_method; + int binary_mode; + int kernel_id; +} Inputs; + typedef struct{ int * num_nucs; // Length = length_num_nucs; double * concs; // Length = length_concs @@ -78,7 +92,7 @@ void binary_write( Inputs in, SimulationData SD ); SimulationData binary_read( Inputs in ); // Simulation.c -unsigned long long run_event_based_simulation(Inputs in, SimulationData SD, int mype, double * kernel_init_time, Profile* profile); +unsigned long long run_event_based_simulation(Inputs in, SimulationData SD, int mype, double * kernel_init_time); int pick_mat(unsigned long * seed); double LCG_random_double(uint64_t * seed); uint64_t fast_forward_LCG(uint64_t seed, uint64_t n); diff --git a/sycl-acc/io.cpp b/sycl-acc/io.cpp index 1ae11f87..460975c6 100644 --- a/sycl-acc/io.cpp +++ b/sycl-acc/io.cpp @@ -86,7 +86,6 @@ int print_results( Inputs in, int mype, double runtime, int nprocs, printf("Lookups/s: "); fancy_int(sim_only_lookups_per_sec); #endif - printf("Total XS Iterations: "); fancy_int(in.num_iterations); } unsigned long long large = 0; @@ -224,7 +223,6 @@ void print_CLI_error(void) printf(" -h Number of hash bins (only relevant when used with \"-G hash\")\n"); printf(" -b Read or write all data structures to file. If reading, this will skip initialization phase. (read, write)\n"); printf(" -k Specifies which kernel to run. 0 is baseline, 1, 2, etc are optimized variants. (0 is default.)\n"); - printf(" -n Specifies how many kernel iterations to run. (1 is default.)\n"); printf("Default is equivalent to: -m history -s large -l 34 -p 500000 -G unionized -n 1\n"); printf("See readme for full description of default run values\n"); exit(4); @@ -264,15 +262,6 @@ Inputs read_CLI( int argc, char * argv[] ) // defaults to baseline kernel input.kernel_id = 0; - // defaults to one kernel iteration - input.num_iterations = 1; - - // default to zero warmup iterations - input.num_warmups = 1; - - // default to stdout - input.filename = NULL; - // defaults to H-M Large benchmark input.HM = (char *) malloc( 6 * sizeof(char) ); input.HM[0] = 'l' ; @@ -412,36 +401,6 @@ Inputs read_CLI( int argc, char * argv[] ) else print_CLI_error(); } - // number of kernel iterations (-n) - else if( strcmp(arg, "-n") == 0 ) - { - if( ++i < argc ) - { - input.num_iterations = atoi(argv[i]); - } - else - print_CLI_error(); - } - else if( strcmp(arg, "--csv") == 0 ) - { - if( ++i < argc ) { - input.filename = (char *)malloc(strlen(argv[i]) + 1); - strcpy(input.filename, argv[i]); - } - else - print_CLI_error(); - } - else if( strcmp(arg, "-w") == 0 ) - { - if( ++i < argc) - { - input.num_warmups = atoi(argv[i]); - } - else - print_CLI_error(); - } - else - print_CLI_error(); } // Validate Input @@ -466,10 +425,6 @@ Inputs read_CLI( int argc, char * argv[] ) if( input.hash_bins < 1 ) print_CLI_error(); - // Validate number of iterations - if( input.num_iterations < 1 ) - print_CLI_error(); - // Validate HM size if( strcasecmp(input.HM, "small") != 0 && strcasecmp(input.HM, "large") != 0 && diff --git a/sycl-usm/Main.cpp b/sycl-usm/Main.cpp index 22d44832..0c475947 100644 --- a/sycl-usm/Main.cpp +++ b/sycl-usm/Main.cpp @@ -68,8 +68,6 @@ int main( int argc, char* argv[] ) border_print(); } - Profile profile; - // Start Simulation Timer omp_start = get_time(); double kernel_init_time; @@ -79,7 +77,7 @@ int main( int argc, char* argv[] ) { if( in.kernel_id == 0 ) { - verification = run_event_based_simulation(in, SD, mype, &kernel_init_time, &profile); + verification = run_event_based_simulation(in, SD, mype, &kernel_init_time); } else { @@ -117,7 +115,5 @@ int main( int argc, char* argv[] ) MPI_Finalize(); #endif - print_profile(profile, in); - return is_invalid_result; } diff --git a/sycl-usm/Simulation.cpp b/sycl-usm/Simulation.cpp index 3227b6d4..2ceea269 100644 --- a/sycl-usm/Simulation.cpp +++ b/sycl-usm/Simulation.cpp @@ -12,7 +12,7 @@ //////////////////////////////////////////////////////////////////////////////////// // use SYCL namespace to reduce symbol names -unsigned long long run_event_based_simulation(Inputs in, SimulationData SD, int mype, double * kernel_init_time, Profile* profile) { +unsigned long long run_event_based_simulation(Inputs in, SimulationData SD, int mype, double * kernel_init_time) { //////////////////////////////////////////////////////////////////////////////// // SUMMARY: Simulation Data Structure Manifest for "SD" Object @@ -51,7 +51,6 @@ unsigned long long run_event_based_simulation(Inputs in, SimulationData SD, int // Create Device Buffers //////////////////////////////////////////////////////////////////////////////// - double startP = get_time(); int* num_nucs = sycl::malloc_device(SD.length_num_nucs, sycl_q); double* concs = sycl::malloc_device(SD.length_concs, sycl_q); int* mats = sycl::malloc_device(SD.length_mats, sycl_q); @@ -77,94 +76,80 @@ unsigned long long run_event_based_simulation(Inputs in, SimulationData SD, int } sycl_q.memcpy(nuclide_grid, SD.nuclide_grid, SD.length_nuclide_grid * sizeof(NuclideGridPoint)); sycl_q.wait(); - profile->host_to_device_time = get_time() - startP; if(mype==0) printf("Beginning event based simulation...\n"); - int nwarmups = in.num_warmups; startP = 0.0; - for (int it = 0; it < in.num_iterations + nwarmups; it++) { - if (it == nwarmups) { - sycl_q.wait(); - startP = get_time(); - } + //////////////////////////////////////////////////////////////////////////////// + // Define Device Kernel + //////////////////////////////////////////////////////////////////////////////// + // queue a kernel to be run, as a lambda + sycl_q.submit([&](sycl::handler &cgh) { //////////////////////////////////////////////////////////////////////////////// - // Define Device Kernel + // XS Lookup Simulation Loop //////////////////////////////////////////////////////////////////////////////// - - // queue a kernel to be run, as a lambda - sycl_q.submit([&](sycl::handler &cgh) { - //////////////////////////////////////////////////////////////////////////////// - // XS Lookup Simulation Loop - //////////////////////////////////////////////////////////////////////////////// - cgh.parallel_for(sycl::range<1>(in.lookups), [=](sycl::id<1> idx) { - // get the index to operate on, first dimemsion - size_t i = idx[0]; - - // Set the initial seed value - uint64_t seed = STARTING_SEED; - - // Forward seed to lookup index (we need 2 samples per lookup) - seed = fast_forward_LCG(seed, 2*i); - - // Randomly pick an energy and material for the particle - double p_energy = LCG_random_double(&seed); - int mat = pick_mat(&seed); - - // debugging - //printf("E = %lf mat = %d\n", p_energy, mat); - - double macro_xs_vector[5] = {0}; - - // Perform macroscopic Cross Section Lookup - calculate_macro_xs( - p_energy, // Sampled neutron energy (in lethargy) - mat, // Sampled material type index neutron is in - in.n_isotopes, // Total number of isotopes in simulation - in.n_gridpoints, // Number of gridpoints per isotope in simulation - num_nucs, // 1-D array with number of nuclides per material - concs, // Flattened 2-D array with concentration of each nuclide in each material - unionized_energy_array, // 1-D Unionized energy array - index_grid, // Flattened 2-D grid holding indices into nuclide grid for each unionized energy level - nuclide_grid, // Flattened 2-D grid holding energy levels and XS_data for all nuclides in simulation - mats, // Flattened 2-D array with nuclide indices defining composition of each type of material - macro_xs_vector, // 1-D array with result of the macroscopic cross section (5 different reaction channels) - in.grid_type, // Lookup type (nuclide, hash, or unionized) - in.hash_bins, // Number of hash bins used (if using hash lookup type) - SD.max_num_nucs // Maximum number of nuclides present in any material - ); - - // For verification, and to prevent the compiler from optimizing - // all work out, we interrogate the returned macro_xs_vector array - // to find its maximum value index, then increment the verification - // value by that index. In this implementation, we store to a global - // array that will get tranferred back and reduced on the host. - double max = -1.0; - int max_idx = 0; - for(int j = 0; j < 5; j++ ) + cgh.parallel_for(sycl::range<1>(in.lookups), [=](sycl::id<1> idx) { + // get the index to operate on, first dimemsion + size_t i = idx[0]; + + // Set the initial seed value + uint64_t seed = STARTING_SEED; + + // Forward seed to lookup index (we need 2 samples per lookup) + seed = fast_forward_LCG(seed, 2*i); + + // Randomly pick an energy and material for the particle + double p_energy = LCG_random_double(&seed); + int mat = pick_mat(&seed); + + // debugging + //printf("E = %lf mat = %d\n", p_energy, mat); + + double macro_xs_vector[5] = {0}; + + // Perform macroscopic Cross Section Lookup + calculate_macro_xs( + p_energy, // Sampled neutron energy (in lethargy) + mat, // Sampled material type index neutron is in + in.n_isotopes, // Total number of isotopes in simulation + in.n_gridpoints, // Number of gridpoints per isotope in simulation + num_nucs, // 1-D array with number of nuclides per material + concs, // Flattened 2-D array with concentration of each nuclide in each material + unionized_energy_array, // 1-D Unionized energy array + index_grid, // Flattened 2-D grid holding indices into nuclide grid for each unionized energy level + nuclide_grid, // Flattened 2-D grid holding energy levels and XS_data for all nuclides in simulation + mats, // Flattened 2-D array with nuclide indices defining composition of each type of material + macro_xs_vector, // 1-D array with result of the macroscopic cross section (5 different reaction channels) + in.grid_type, // Lookup type (nuclide, hash, or unionized) + in.hash_bins, // Number of hash bins used (if using hash lookup type) + SD.max_num_nucs // Maximum number of nuclides present in any material + ); + + // For verification, and to prevent the compiler from optimizing + // all work out, we interrogate the returned macro_xs_vector array + // to find its maximum value index, then increment the verification + // value by that index. In this implementation, we store to a global + // array that will get tranferred back and reduced on the host. + double max = -1.0; + int max_idx = 0; + for(int j = 0; j < 5; j++ ) + { + if( macro_xs_vector[j] > max ) { - if( macro_xs_vector[j] > max ) - { - max = macro_xs_vector[j]; - max_idx = j; - } + max = macro_xs_vector[j]; + max_idx = j; } - verification[i] = max_idx+1; + } + verification[i] = max_idx+1; - }); }); - } + }); stop = get_time(); - sycl_q.wait(); - profile->kernel_time = get_time() - startP; - if(mype==0) printf("Kernel initialization, compilation, and launch took %.2lf seconds.\n", stop-start); - startP = get_time(); sycl_q.memcpy(verification_host, verification, in.lookups * sizeof(int)).wait(); - profile->device_to_host_time = get_time() - startP; // Host reduces the verification array unsigned long long verification_scalar = 0; diff --git a/sycl-usm/XSbench_header.h b/sycl-usm/XSbench_header.h index 5e492ee1..2f00da3b 100644 --- a/sycl-usm/XSbench_header.h +++ b/sycl-usm/XSbench_header.h @@ -11,7 +11,6 @@ #include #include #include -#include "../XSbench_shared_header.h" // Papi Header #ifdef PAPI @@ -45,6 +44,20 @@ typedef struct{ double nu_fission_xs; } NuclideGridPoint; +typedef struct{ + int nthreads; + long n_isotopes; + long n_gridpoints; + int lookups; + char * HM; + int grid_type; // 0: Unionized Grid (default) 1: Nuclide Grid + int hash_bins; + int particles; + int simulation_method; + int binary_mode; + int kernel_id; +} Inputs; + typedef struct{ int * num_nucs; // Length = length_num_nucs; double * concs; // Length = length_concs @@ -78,7 +91,7 @@ void binary_write( Inputs in, SimulationData SD ); SimulationData binary_read( Inputs in ); // Simulation.c -unsigned long long run_event_based_simulation(Inputs in, SimulationData SD, int mype, double * kernel_init_time, Profile* profile); +unsigned long long run_event_based_simulation(Inputs in, SimulationData SD, int mype, double * kernel_init_time); int pick_mat(unsigned long * seed); double LCG_random_double(uint64_t * seed); uint64_t fast_forward_LCG(uint64_t seed, uint64_t n); diff --git a/sycl-usm/io.cpp b/sycl-usm/io.cpp index 1ae11f87..460975c6 100644 --- a/sycl-usm/io.cpp +++ b/sycl-usm/io.cpp @@ -86,7 +86,6 @@ int print_results( Inputs in, int mype, double runtime, int nprocs, printf("Lookups/s: "); fancy_int(sim_only_lookups_per_sec); #endif - printf("Total XS Iterations: "); fancy_int(in.num_iterations); } unsigned long long large = 0; @@ -224,7 +223,6 @@ void print_CLI_error(void) printf(" -h Number of hash bins (only relevant when used with \"-G hash\")\n"); printf(" -b Read or write all data structures to file. If reading, this will skip initialization phase. (read, write)\n"); printf(" -k Specifies which kernel to run. 0 is baseline, 1, 2, etc are optimized variants. (0 is default.)\n"); - printf(" -n Specifies how many kernel iterations to run. (1 is default.)\n"); printf("Default is equivalent to: -m history -s large -l 34 -p 500000 -G unionized -n 1\n"); printf("See readme for full description of default run values\n"); exit(4); @@ -264,15 +262,6 @@ Inputs read_CLI( int argc, char * argv[] ) // defaults to baseline kernel input.kernel_id = 0; - // defaults to one kernel iteration - input.num_iterations = 1; - - // default to zero warmup iterations - input.num_warmups = 1; - - // default to stdout - input.filename = NULL; - // defaults to H-M Large benchmark input.HM = (char *) malloc( 6 * sizeof(char) ); input.HM[0] = 'l' ; @@ -412,36 +401,6 @@ Inputs read_CLI( int argc, char * argv[] ) else print_CLI_error(); } - // number of kernel iterations (-n) - else if( strcmp(arg, "-n") == 0 ) - { - if( ++i < argc ) - { - input.num_iterations = atoi(argv[i]); - } - else - print_CLI_error(); - } - else if( strcmp(arg, "--csv") == 0 ) - { - if( ++i < argc ) { - input.filename = (char *)malloc(strlen(argv[i]) + 1); - strcpy(input.filename, argv[i]); - } - else - print_CLI_error(); - } - else if( strcmp(arg, "-w") == 0 ) - { - if( ++i < argc) - { - input.num_warmups = atoi(argv[i]); - } - else - print_CLI_error(); - } - else - print_CLI_error(); } // Validate Input @@ -466,10 +425,6 @@ Inputs read_CLI( int argc, char * argv[] ) if( input.hash_bins < 1 ) print_CLI_error(); - // Validate number of iterations - if( input.num_iterations < 1 ) - print_CLI_error(); - // Validate HM size if( strcasecmp(input.HM, "small") != 0 && strcasecmp(input.HM, "large") != 0 &&