Skip to content

Commit 4f34313

Browse files
tobre1tobre1
committed
Add flux testing model
1 parent fd7c89c commit 4f34313

6 files changed

Lines changed: 225 additions & 0 deletions

File tree

‎gpu/include/models/CallableWrapper.cu‎

Lines changed: 9 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -6,6 +6,7 @@
66
#include "PlasmaEtching.cuh"
77
#include "SingleParticle.cuh"
88
#include "SingleParticleALD.cuh"
9+
#include "TestFlux.cuh"
910

1011
//
1112
// --- Direct Callables wrapper
@@ -147,3 +148,11 @@ __direct_callable__faradayCollision(const void *sbtData,
147148
viennaray::gpu::PerRayData *prd) {
148149
faradayIonCollision(sbtData, prd);
149150
}
151+
152+
// --- TestFlux particle
153+
154+
extern "C" __device__ void
155+
__direct_callable__testFluxCollision(const void *,
156+
viennaray::gpu::PerRayData *prd) {
157+
testFluxCollision(prd);
158+
}

‎gpu/include/models/TestFlux.cuh‎

Lines changed: 32 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,32 @@
1+
#pragma once
2+
3+
#include <vcContext.hpp>
4+
#include <vcVectorType.hpp>
5+
6+
#include <raygLaunchParams.hpp>
7+
#include <raygReflection.hpp>
8+
9+
extern "C" __constant__ viennaray::gpu::LaunchParams launchParams;
10+
11+
//
12+
// --- TestFlux particle
13+
//
14+
15+
__forceinline__ __device__ void
16+
testFluxCollision(viennaray::gpu::PerRayData *prd) {
17+
for (int i = 0; i < prd->ISCount; ++i) {
18+
if (prd->numReflections == 0) {
19+
atomicAdd(
20+
&launchParams
21+
.resultBuffer[viennaray::gpu::getIdxOffset(0, launchParams) +
22+
prd->primIDs[i]],
23+
prd->rayWeight);
24+
} else {
25+
atomicAdd(
26+
&launchParams
27+
.resultBuffer[viennaray::gpu::getIdxOffset(1, launchParams) +
28+
prd->primIDs[i]],
29+
prd->rayWeight);
30+
}
31+
}
32+
}

‎include/viennaps/models/psTestFluxModel.hpp‎

Lines changed: 138 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,138 @@
1+
#pragma once
2+
3+
#include "../process/psProcessModel.hpp"
4+
#include "../psMaterials.hpp"
5+
6+
#include <rayParticle.hpp>
7+
#include <rayReflection.hpp>
8+
9+
namespace viennaps {
10+
11+
using namespace viennacore;
12+
13+
namespace impl {
14+
template <typename NumericType, int D>
15+
class TestFluxParticle
16+
: public viennaray::Particle<TestFluxParticle<NumericType, D>,
17+
NumericType> {
18+
public:
19+
TestFluxParticle(NumericType sticking, NumericType sourcePower)
20+
: stickingProbability_(sticking), sourcePower_(sourcePower) {}
21+
22+
void surfaceCollision(NumericType rayWeight, const Vec3D<NumericType> &,
23+
const Vec3D<NumericType> &, const unsigned int primID,
24+
const int,
25+
viennaray::TracingData<NumericType> &localData,
26+
const viennaray::TracingData<NumericType> *,
27+
RNG &) override final {
28+
if (!isReflected_) {
29+
localData.getVectorData(0)[primID] += rayWeight;
30+
} else {
31+
localData.getVectorData(1)[primID] += rayWeight;
32+
}
33+
}
34+
std::pair<NumericType, Vec3D<NumericType>>
35+
surfaceReflection(NumericType, const Vec3D<NumericType> &,
36+
const Vec3D<NumericType> &geomNormal, const unsigned int,
37+
const int, const viennaray::TracingData<NumericType> *,
38+
RNG &rngState) override final {
39+
isReflected_ = true;
40+
auto direction =
41+
viennaray::ReflectionDiffuse<NumericType, D>(geomNormal, rngState);
42+
return std::pair<NumericType, Vec3D<NumericType>>{stickingProbability_,
43+
direction};
44+
}
45+
void initNew(RNG &) override final { isReflected_ = false; }
46+
NumericType getSourceDistributionPower() const override final {
47+
return sourcePower_;
48+
}
49+
std::vector<std::string> getLocalDataLabels() const override final {
50+
return {"primaryFlux", "secondaryFlux"};
51+
}
52+
53+
private:
54+
const NumericType stickingProbability_;
55+
const NumericType sourcePower_;
56+
bool isReflected_ = false;
57+
};
58+
} // namespace impl
59+
60+
#ifdef VIENNACORE_COMPILE_GPU
61+
namespace gpu {
62+
63+
template <typename NumericType, int D>
64+
class TestFluxModel final : public ProcessModelGPU<NumericType, D> {
65+
public:
66+
TestFluxModel(NumericType stickingProbability, NumericType sourceExponent) {
67+
// particles
68+
viennaray::gpu::Particle<NumericType> particle{
69+
.name = "TestFluxParticle",
70+
.sticking = stickingProbability,
71+
.cosineExponent = sourceExponent};
72+
particle.dataLabels.push_back("primaryFlux");
73+
particle.dataLabels.push_back("secondaryFlux");
74+
75+
std::unordered_map<std::string, unsigned> pMap = {{"TestFluxParticle", 0}};
76+
std::vector<viennaray::gpu::CallableConfig> cMap = {
77+
{0, viennaray::gpu::CallableSlot::COLLISION,
78+
"__direct_callable__testFluxCollision"},
79+
{0, viennaray::gpu::CallableSlot::REFLECTION,
80+
"__direct_callable__singleNeutralReflection"}};
81+
this->setParticleCallableMap(pMap, cMap);
82+
this->setCallableFileName("CallableWrapper");
83+
84+
// surface model
85+
auto surfModel = SmartPointer<::viennaps::SurfaceModel<NumericType>>::New();
86+
87+
// velocity field
88+
auto velField =
89+
SmartPointer<::viennaps::DefaultVelocityField<NumericType, D>>::New();
90+
91+
this->setSurfaceModel(surfModel);
92+
this->setVelocityField(velField);
93+
this->insertNextParticleType(particle);
94+
this->setProcessName("TestFluxModel");
95+
this->hasGPU = true;
96+
}
97+
};
98+
} // namespace gpu
99+
#endif
100+
101+
// Etching or deposition based on a single particle model with diffuse
102+
// reflections.
103+
template <typename NumericType, int D>
104+
class TestFluxModel : public ProcessModelCPU<NumericType, D> {
105+
NumericType stickingProbability_ = 1.;
106+
NumericType sourceDistributionPower_ = 1.;
107+
108+
public:
109+
TestFluxModel(NumericType stickingProbability = 1.,
110+
NumericType sourceDistributionPower = 1.)
111+
: stickingProbability_(stickingProbability),
112+
sourceDistributionPower_(sourceDistributionPower) {
113+
114+
auto particle = std::make_unique<impl::TestFluxParticle<NumericType, D>>(
115+
stickingProbability_, sourceDistributionPower_);
116+
117+
// surface model
118+
auto surfModel = SmartPointer<SurfaceModel<NumericType>>::New();
119+
120+
// velocity field
121+
auto velField = SmartPointer<DefaultVelocityField<NumericType, D>>::New();
122+
123+
this->setSurfaceModel(surfModel);
124+
this->setVelocityField(velField);
125+
this->insertNextParticleType(particle);
126+
this->setProcessName("TestFluxModel");
127+
this->hasGPU = true;
128+
}
129+
130+
#ifdef VIENNACORE_COMPILE_GPU
131+
SmartPointer<ProcessModelBase<NumericType, D>> getGPUModel() final {
132+
return SmartPointer<gpu::TestFluxModel<NumericType, D>>::New(
133+
stickingProbability_, sourceDistributionPower_);
134+
}
135+
#endif
136+
};
137+
138+
} // namespace viennaps

‎python/pyWrapDimension.hpp‎

Lines changed: 13 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -68,6 +68,7 @@
6868
#include <models/psSingleParticleProcess.hpp>
6969
#include <models/psTEOSDeposition.hpp>
7070
#include <models/psTEOSPECVD.hpp>
71+
#include <models/psTestFluxModel.hpp>
7172
#include <models/psWetEtching.hpp>
7273

7374
// visualization
@@ -762,6 +763,12 @@ template <int D> void bindApi(py::module &module) {
762763
py::arg("label") = "ionFlux")
763764
.def("setRateFunction", &MultiParticleProcess<T, D>::setRateFunction);
764765

766+
// Test Flux Model
767+
py::class_<TestFluxModel<T, D>, SmartPointer<TestFluxModel<T, D>>>(
768+
module, "TestFluxModel", processModel)
769+
.def(py::init<T, T>(), py::arg("stickingProbability"),
770+
py::arg("sourceExponent"));
771+
765772
// TEOS Deposition
766773
py::class_<TEOSDeposition<T, D>, SmartPointer<TEOSDeposition<T, D>>>(
767774
module, "TEOSDeposition", processModel)
@@ -1463,6 +1470,12 @@ template <int D> void bindApi(py::module &module) {
14631470
.def("setRateFunction",
14641471
&gpu::MultiParticleProcess<T, D>::setRateFunction);
14651472

1473+
// Test Flux Model
1474+
py::class_<gpu::TestFluxModel<T, D>, SmartPointer<gpu::TestFluxModel<T, D>>>(
1475+
m_gpu, "TestFluxModel", processModel_gpu)
1476+
.def(py::init<T, T>(), py::arg("stickingProbability"),
1477+
py::arg("sourceExponent"));
1478+
14661479
// SF6O2 Etching
14671480
py::class_<gpu::SF6O2Etching<T, D>, SmartPointer<gpu::SF6O2Etching<T, D>>>(
14681481
m_gpu, "SF6O2Etching", processModel_gpu)

‎tests/testFluxModel/CMakeLists.txt‎

Lines changed: 7 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,7 @@
1+
project(testFluxModel LANGUAGES CXX)
2+
3+
add_gpu_executable(${PROJECT_NAME} target_name "${PROJECT_NAME}.cpp")
4+
# target_link_libraries(${PROJECT_NAME} PRIVATE ViennaPS)
5+
6+
add_dependencies(ViennaPS_Tests ${PROJECT_NAME})
7+
add_test(NAME ${PROJECT_NAME} COMMAND $<TARGET_FILE:${PROJECT_NAME}>)

‎tests/testFluxModel/testFluxModel.cpp‎

Lines changed: 26 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,26 @@
1+
#include <geometries/psMakeHole.hpp>
2+
#include <models/psTestFluxModel.hpp>
3+
#include <process/psProcess.hpp>
4+
5+
int main() {
6+
using NumericType = double;
7+
constexpr int D = 3;
8+
9+
// Create domain
10+
auto domain = viennaps::Domain<NumericType, D>::New(1.0, 20., 20.);
11+
viennaps::MakeHole<NumericType, D>(domain, 8.0, 10.0).apply();
12+
13+
// Create process model
14+
auto processModel =
15+
viennaps::SmartPointer<viennaps::TestFluxModel<NumericType, D>>::New(0.3,
16+
1.0);
17+
18+
// Create process
19+
viennaps::Process<NumericType, D> process(domain, processModel, 1.0);
20+
process.setFluxEngineType(viennaps::FluxEngineType::GPU_TRIANGLE);
21+
22+
// Execute process
23+
auto mesh = process.calculateFlux();
24+
25+
return 0;
26+
}

0 commit comments

Comments
 (0)