Skip to content

Commit 3fd2d6e

Browse files
Vedant2005goyalVedant2005goyal
authored
Add XSBench Benchmark (#1878)
1 parent 783760d commit 3fd2d6e

2 files changed

Lines changed: 136 additions & 0 deletions

File tree

‎benchmark/CMakeLists.txt‎

Lines changed: 31 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -102,6 +102,37 @@ if(CUDAToolkit_FOUND)
102102
)
103103

104104
target_include_directories(LULESH_Benchmark PRIVATE ${lulesh-cuda_SOURCE_DIR}/LULESH)
105+
execute_process(
106+
COMMAND git apply --check "${CMAKE_CURRENT_SOURCE_DIR}/XSBench_diff.patch"
107+
WORKING_DIRECTORY "${lulesh-cuda_SOURCE_DIR}/XSBench"
108+
RESULT_VARIABLE xsbench_patch_applied_or_not
109+
)
110+
111+
if(xsbench_patch_applied_or_not EQUAL 0)
112+
execute_process(
113+
COMMAND git apply "${CMAKE_CURRENT_SOURCE_DIR}/XSBench_diff.patch"
114+
WORKING_DIRECTORY "${lulesh-cuda_SOURCE_DIR}/XSBench"
115+
)
116+
endif()
117+
118+
set(XSBENCH_SOURCES
119+
${lulesh-cuda_SOURCE_DIR}/XSBench/Main.cu
120+
${lulesh-cuda_SOURCE_DIR}/XSBench/io.cu
121+
${lulesh-cuda_SOURCE_DIR}/XSBench/Simulation.cu
122+
${lulesh-cuda_SOURCE_DIR}/XSBench/GridInit.cu
123+
${lulesh-cuda_SOURCE_DIR}/XSBench/XSutils.cu
124+
${lulesh-cuda_SOURCE_DIR}/XSBench/Materials.cu
125+
)
126+
127+
set_source_files_properties(${XSBENCH_SOURCES} PROPERTIES LANGUAGE CXX)
128+
CB_ADD_CUDA_BENCHMARK(XSBench_Benchmark ${XSBENCH_SOURCES})
129+
target_compile_options(XSBench_Benchmark PRIVATE -fno-openmp)
130+
target_compile_definitions(XSBench_Benchmark PRIVATE
131+
ALWAYS_INLINE=1
132+
USE_CLAD=1
133+
)
134+
135+
target_include_directories(XSBench_Benchmark PRIVATE ${lulesh-cuda_SOURCE_DIR}/XSBench)
105136
endif()
106137

107138

‎benchmark/XSBench_diff.patch‎

Lines changed: 105 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,105 @@
1+
diff --git a/XSBench/Simulation.cu b/XSBench/Simulation.cu
2+
index 59ffad6..3f7099d 100644
3+
--- a/XSBench/Simulation.cu
4+
+++ b/XSBench/Simulation.cu
5+
@@ -1,5 +1,7 @@
6+
#include "XSbench_header.cuh"
7+
-
8+
+#ifdef USE_CLAD
9+
+#include "clad/Differentiator/Differentiator.h"
10+
+#endif
11+
////////////////////////////////////////////////////////////////////////////////////
12+
// BASELINE FUNCTIONS
13+
////////////////////////////////////////////////////////////////////////////////////
14+
@@ -54,7 +56,7 @@ unsigned long long run_event_based_simulation_baseline(Inputs in, SimulationData
15+
16+
return verification_scalar;
17+
}
18+
-
19+
+#ifndef USE_CLAD
20+
template<typename... Args>
21+
__device__
22+
void __enzyme_autodiff(void*, Args...);
23+
@@ -64,6 +66,7 @@ __device__ int enzyme_dup, enzyme_const, enzyme_active;
24+
// In this kernel, we perform a single lookup with each thread. Threads within a warp
25+
// do not really have any relation to each other, and divergence due to high nuclide count fuel
26+
// material lookups are costly. This kernel constitutes baseline performance.
27+
+#endif
28+
__global__ void xs_lookup_kernel_baseline(Inputs in, SimulationData GSD )
29+
{
30+
// The lookup ID. Used to set the seed, and to store the verification value
31+
@@ -133,6 +136,16 @@ __global__ void xs_lookup_kernel_baseline(Inputs in, SimulationData GSD )
32+
// printf("i=%d dd=%f, out2=%f out1=%f | in2=%f in1=%f\n", i, (macro_xs_vector2[0] - macro_xs_vector[0]) / DELTA, macro_xs_vector2[0], macro_xs_vector[0], GSD.d_nuclide_grid[0].energy, GSD.nuclide_grid[0].energy);
33+
atomicAdd(GSD.dout, (macro_xs_vector2[0] - macro_xs_vector[0]) / DELTA );
34+
#endif
35+
+ #else
36+
+ #ifdef USE_CLAD
37+
+ clad::gradient(calculate_macro_xs, "nuclide_grids, macro_xs_vector")
38+
+ .execute(
39+
+ p_energy, mat, in.n_isotopes, in.n_gridpoints,
40+
+ GSD.num_nucs, GSD.concs, GSD.unionized_energy_array,
41+
+ GSD.index_grid, GSD.nuclide_grid, GSD.mats,
42+
+ macro_xs_vector, in.grid_type, in.hash_bins, GSD.max_num_nucs,
43+
+ GSD.d_nuclide_grid, d_macro_xs_vector
44+
+ );
45+
#else
46+
__enzyme_autodiff((void*)calculate_macro_xs,
47+
enzyme_const, p_energy, // Sampled neutron energy (in lethargy)
48+
@@ -156,6 +169,7 @@ __global__ void xs_lookup_kernel_baseline(Inputs in, SimulationData GSD )
49+
enzyme_const, in.hash_bins, // Number of hash bins used (if using hash lookup type)
50+
enzyme_const, GSD.max_num_nucs // Maximum number of nuclides present in any material
51+
);
52+
+ #endif
53+
#endif
54+
55+
// For verification, and to prevent the compiler from optimizing
56+
diff --git a/XSBench/XSbench_header.cuh b/XSBench/XSbench_header.cuh
57+
index f0d8639..75afd2c 100644
58+
--- a/XSBench/XSbench_header.cuh
59+
+++ b/XSBench/XSbench_header.cuh
60+
@@ -8,8 +8,19 @@
61+
#include<assert.h>
62+
#include<cuda.h>
63+
#include <thrust/reduce.h>
64+
+#include <thrust/count.h>
65+
+#include <thrust/sort.h>
66+
+#include <thrust/partition.h>
67+
#include<stdint.h>
68+
#include <chrono>
69+
+#ifdef USE_CLAD
70+
+// Marks mutable pointer and array parameters that are not differentiation
71+
+// targets as non-differentiable to avoid Clad's "dependent non-const pointer"
72+
+// error without changing XSBench's function signatures.
73+
+#define CLAD_NON_DIFF __attribute__((annotate("non_differentiable")))
74+
+#else
75+
+#define CLAD_NON_DIFF
76+
+#endif
77+
78+
// Grid types
79+
#define UNIONIZED 0
80+
@@ -107,18 +118,18 @@ unsigned long long run_event_based_simulation_baseline(Inputs in, SimulationData
81+
__global__ void xs_lookup_kernel_baseline(Inputs in, SimulationData GSD );
82+
__device__ void calculate_micro_xs( double p_energy, int nuc, long n_isotopes,
83+
long n_gridpoints,
84+
- double * __restrict__ egrid, int * __restrict__ index_data,
85+
+ CLAD_NON_DIFF double * __restrict__ egrid, CLAD_NON_DIFF int * __restrict__ index_data,
86+
NuclideGridPoint * __restrict__ nuclide_grids,
87+
long idx, double * __restrict__ xs_vector, int grid_type, int hash_bins );
88+
__device__ void calculate_macro_xs( double p_energy, int mat, long n_isotopes,
89+
- long n_gridpoints, int * __restrict__ num_nucs,
90+
- double * __restrict__ concs,
91+
- double * __restrict__ egrid, int * __restrict__ index_data,
92+
+ long n_gridpoints, CLAD_NON_DIFF int * __restrict__ num_nucs,
93+
+ CLAD_NON_DIFF double * __restrict__ concs,
94+
+ CLAD_NON_DIFF double * __restrict__ egrid, CLAD_NON_DIFF int * __restrict__ index_data,
95+
NuclideGridPoint * __restrict__ nuclide_grids,
96+
- int * __restrict__ mats,
97+
+ CLAD_NON_DIFF int * __restrict__ mats,
98+
double * __restrict__ macro_xs_vector, int grid_type, int hash_bins, int max_num_nucs );
99+
-__host__ __device__ long grid_search( long n, double quarry, double * __restrict__ A);
100+
-__host__ __device__ long grid_search_nuclide( long n, double quarry, NuclideGridPoint * A, long low, long high);
101+
+CLAD_NON_DIFF __host__ __device__ long grid_search( long n, double quarry, double * __restrict__ A);
102+
+CLAD_NON_DIFF __host__ __device__ long grid_search_nuclide( long n, double quarry, NuclideGridPoint * A, long low, long high);
103+
__device__ int pick_mat( uint64_t * seed );
104+
__host__ __device__ double LCG_random_double(uint64_t * seed);
105+
__device__ uint64_t fast_forward_LCG(uint64_t seed, uint64_t n);

0 commit comments

Comments
 (0)