Skip to content

Upstreaming #23

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Open
joy-kitson wants to merge 5 commits into
base: master
Choose a base branch
from
Open

Upstreaming #23

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
5 commits
Select commit Hold shift + click to select a range
2cc6093
style: moved Inputs struct to individual header files and removed ext…
joy-kitson Aug 18, 2026
b738c2d
style: removed references to the profile object and additional timer …
joy-kitson Aug 19, 2026
9237c02
style: removed reference to warmups and multiple iterations
joy-kitson Aug 19, 2026
ce45c45
fix: updated Kokkos CMakeLists.txt to properly check whether CUDA or …
joy-kitson Sep 14, 2026
a0923bd
fix: now properly checking for cuda and hip with raja
joy-kitson Sep 15, 2026
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
17 changes: 0 additions & 17 deletions XSbench_shared_header.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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;
Expand Down
6 changes: 1 addition & 5 deletions cuda/Main.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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
Expand All @@ -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)
Expand Down Expand Up @@ -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;
}
16 changes: 2 additions & 14 deletions cuda/Simulation.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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
Expand All @@ -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<<<nblocks, nthreads>>>( in, GSD );
}
xs_lookup_kernel_baseline<<<nblocks, nthreads>>>( 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++ )
Expand Down
17 changes: 15 additions & 2 deletions cuda/XSbench_header.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -10,7 +10,6 @@
#include <thrust/partition.h>
#include<stdint.h>
#include <chrono>
#include "../XSbench_shared_header.h"

// Grid types
#define UNIONIZED 0
Expand Down Expand Up @@ -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
Expand Down Expand Up @@ -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,
Expand Down
47 changes: 0 additions & 47 deletions cuda/io.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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);
Expand Down Expand Up @@ -219,9 +218,6 @@ void print_CLI_error(void)
printf(" -h <hash bins> Number of hash bins (only relevant when used with \"-G hash\")\n");
printf(" -b <binary mode> Read or write all data structures to file. If reading, this will skip initialization phase. (read, write)\n");
printf(" -k <kernel ID> Specifies which kernel to run. 0 is baseline, 1, 2, etc are optimized variants. (0 is default.)\n");
printf(" -n <num iterations> Specifies how many kernel iterations to run. (1 is default.)\n");
printf(" -w <num warmups> Specifies how many warmup iterations to run. (0 is default.)\n");
printf(" --csv <file path> 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);
Expand Down Expand Up @@ -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' ;
Expand Down Expand Up @@ -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
Expand All @@ -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 &&
Expand Down
8 changes: 1 addition & 7 deletions hip/Main.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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
Expand All @@ -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);
Expand Down Expand Up @@ -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;
}
15 changes: 2 additions & 13 deletions hip/Simulation.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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
Expand All @@ -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
Expand Down
17 changes: 15 additions & 2 deletions hip/XSbench_header.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -9,7 +9,6 @@
#include<hip/hip_runtime.h>
#include<stdint.h>
#include <chrono>
#include "../XSbench_shared_header.h"

// Grid types
#define UNIONIZED 0
Expand Down Expand Up @@ -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
Expand Down Expand Up @@ -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,
Expand Down
45 changes: 0 additions & 45 deletions hip/io.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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);
Expand Down Expand Up @@ -219,7 +218,6 @@ void print_CLI_error(void)
printf(" -h <hash bins> Number of hash bins (only relevant when used with \"-G hash\")\n");
printf(" -b <binary mode> Read or write all data structures to file. If reading, this will skip initialization phase. (read, write)\n");
printf(" -k <kernel ID> Specifies which kernel to run. 0 is baseline, 1, 2, etc are optimized variants. (0 is default.)\n");
printf(" -n <num. iterations> 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);
Expand Down Expand Up @@ -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' ;
Expand Down Expand Up @@ -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
Expand All @@ -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 &&
Expand Down
Loading
Loading