forked from celeritas-project/celeritas
-
Notifications
You must be signed in to change notification settings - Fork 0
/
Physics.test.cu
59 lines (50 loc) · 2.07 KB
/
Physics.test.cu
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
//---------------------------------*-CUDA-*----------------------------------//
// Copyright 2021-2024 UT-Battelle, LLC, and other Celeritas developers.
// See the top-level COPYRIGHT file for details.
// SPDX-License-Identifier: (Apache-2.0 OR MIT)
//---------------------------------------------------------------------------//
//! \file celeritas/phys/Physics.test.cu
//---------------------------------------------------------------------------//
#include "Physics.test.hh"
#include "corecel/device_runtime_api.h"
#include "corecel/sys/Device.hh"
#include "corecel/sys/KernelParamCalculator.device.hh"
#include "celeritas/Quantities.hh"
#include "celeritas/phys/PhysicsStepView.hh"
#include "celeritas/phys/PhysicsTrackView.hh"
namespace celeritas
{
namespace test
{
//---------------------------------------------------------------------------//
namespace
{
//---------------------------------------------------------------------------//
// KERNELS
//---------------------------------------------------------------------------//
__global__ void phys_test_kernel(PTestInput const inp)
{
auto tid = TrackSlotId{KernelParamCalculator::thread_id().unchecked_get()};
if (tid.get() >= inp.states.size())
return;
auto const& init = inp.inits[tid];
PhysicsTrackView phys(inp.params, inp.states, init.particle, init.mat, tid);
PhysicsStepView step(inp.params, inp.states, tid);
phys = PhysicsTrackInitializer{};
inp.result[tid.get()]
= native_value_to<units::CmLength>(calc_step(phys, step, init.energy))
.value();
}
} // namespace
//---------------------------------------------------------------------------//
// TESTING INTERFACE
//---------------------------------------------------------------------------//
//! Run on device and return results
void phys_cuda_test(PTestInput const& input)
{
CELER_ASSERT(input.inits.size() == input.states.size());
CELER_LAUNCH_KERNEL(phys_test, input.states.size(), 0, input);
}
//---------------------------------------------------------------------------//
} // namespace test
} // namespace celeritas