Skip to content

Commit 59e8e2a

Browse files
tobre1tobre1
authored
fix: HBrO2 model configuration and material handling on GPU (#194)
* fix: material mapping on GPU, threshold energies for HBrO2 model * fix: add dependencies to nvcc compilation to trigger recompilation * fix: DRAM wiggling configuration
1 parent 589823f commit 59e8e2a

10 files changed

Lines changed: 117 additions & 105 deletions

File tree

‎CMakeLists.txt‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -105,7 +105,7 @@ include("cmake/cpm.cmake")
105105

106106
CPMAddPackage(
107107
NAME ViennaCore
108-
VERSION 1.9.1
108+
VERSION 1.9.2
109109
GIT_REPOSITORY "https://github.com/ViennaTools/ViennaCore"
110110
EXCLUDE_FROM_ALL ${VIENNAPS_BUILD_PYTHON}
111111
OPTIONS "VIENNACORE_USE_GPU ${VIENNAPS_USE_GPU}")
@@ -117,7 +117,7 @@ CPMAddPackage(
117117

118118
CPMAddPackage(
119119
NAME ViennaRay
120-
VERSION 3.10.0
120+
VERSION 3.10.1
121121
GIT_REPOSITORY "https://github.com/ViennaTools/ViennaRay"
122122
EXCLUDE_FROM_ALL ${VIENNAPS_BUILD_PYTHON}
123123
OPTIONS "VIENNARAY_USE_GPU ${VIENNAPS_USE_GPU}")

‎examples/DRAMWiggling/DRAMWiggling.cpp‎

Lines changed: 1 addition & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -6,9 +6,6 @@ int main(int argc, char **argv) {
66
using NumericType = double;
77
constexpr int D = 3;
88

9-
Logger::setLogLevel(LogLevel::ERROR);
10-
omp_set_num_threads(12);
11-
129
// Parse the parameters
1310
util::Parameters params;
1411
if (argc > 1) {
@@ -51,6 +48,7 @@ int main(int argc, char **argv) {
5148
modelParams.Ions.sigmaEnergy = params.get("sigmaEnergy");
5249
modelParams.Ions.exponent = params.get("ionExponent");
5350
modelParams.Ions.n_l = 200;
51+
modelParams.Substrate.B_sp = 0.75;
5452
auto model = SmartPointer<HBrO2Etching<NumericType, D>>::New(modelParams);
5553

5654
// Advection parameters

‎examples/DRAMWiggling/DRAMWiggling.py‎

Lines changed: 11 additions & 14 deletions
Original file line numberDiff line numberDiff line change
@@ -1,4 +1,3 @@
1-
import viennaps.d3 as psd
21
import viennaps as ps
32
from argparse import ArgumentParser
43

@@ -10,6 +9,8 @@
109
parser.add_argument("filename")
1110
args = parser.parse_args()
1211

12+
ps.setDimension(3)
13+
1314
gridDelta = 0.01 * (1.0 + 1e-12)
1415
boundaryConds = [
1516
ps.BoundaryType.REFLECTIVE_BOUNDARY,
@@ -19,36 +20,34 @@
1920

2021
params = ps.readConfigFile(args.filename)
2122

22-
mask = psd.GDSGeometry(gridDelta, boundaryConds)
23+
mask = ps.GDSGeometry(gridDelta, boundaryConds)
2324
mask.setBoundaryPadding(0.1, 0.1)
24-
reader = psd.GDSReader(mask, params["gdsFile"])
25+
reader = ps.GDSReader(mask, params["gdsFile"])
2526
reader.apply()
2627

2728
# Prepare geometry
28-
geometry = psd.Domain()
29+
geometry = ps.Domain()
2930

3031
# Insert GDS layers
3132
maskLS = mask.layerToLevelSet(0, 0.0, 0.18)
3233
geometry.insertNextLevelSetAsMaterial(maskLS, ps.Material.Mask)
3334

3435
# Add plane
35-
psd.MakePlane(geometry, 0.0, ps.Material.Si, True).apply()
36-
37-
# print intermediate output surfaces during the process
38-
ps.Logger.setLogLevel(ps.LogLevel.INFO)
36+
ps.MakePlane(geometry, 0.0, ps.Material.Si, True).apply()
3937

4038
ps.Length.setUnit(params["lengthUnit"])
4139
ps.Time.setUnit(params["timeUnit"])
4240

43-
modelParams = psd.HBrO2Etching.defaultParameters()
41+
modelParams = ps.HBrO2Etching.defaultParameters()
4442
modelParams.ionFlux = params["ionFlux"]
4543
modelParams.etchantFlux = params["etchantFlux"]
4644
modelParams.passivationFlux = params["oxygenFlux"]
4745
modelParams.Ions.meanEnergy = params["meanEnergy"]
4846
modelParams.Ions.sigmaEnergy = params["sigmaEnergy"]
4947
modelParams.Ions.exponent = params["ionExponent"]
5048
modelParams.Ions.n_l = 200
51-
model = psd.HBrO2Etching(modelParams)
49+
modelParams.Substrate.B_sp = 0.75
50+
model = ps.HBrO2Etching(modelParams)
5251

5352
coverageParameters = ps.CoverageParameters()
5453
coverageParameters.maxIterations = 10
@@ -57,12 +56,10 @@
5756
rayTracingParams.raysPerPoint = int(params["raysPerPoint"])
5857

5958
advectionParams = ps.AdvectionParameters()
60-
advectionParams.spatialScheme = ps.util.convertSpatialScheme(
61-
params["spatialScheme"]
62-
)
59+
advectionParams.spatialScheme = ps.util.convertSpatialScheme(params["spatialScheme"])
6360

6461
# process setup
65-
process = psd.Process(geometry, model)
62+
process = ps.Process(geometry, model)
6663
process.setProcessDuration(params["processTime"]) # seconds
6764
process.setParameters(coverageParameters)
6865
process.setParameters(rayTracingParams)

‎examples/DRAMWiggling/config.txt‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -19,6 +19,6 @@ sigmaEnergy=10 # eV
1919

2020
spatialScheme=LF_2
2121

22-
numSteps=100
22+
numSteps=20
2323
raysPerPoint=1000
2424
gdsFile=wiggle_full.gds

‎gpu/CMakeLists.txt‎

Lines changed: 11 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -11,7 +11,17 @@ if(NOT VIENNAPS_PASSED_FIRST_CONFIGURE)
1111
endif()
1212

1313
# Add the callable wrapper
14-
viennacore_add_optixir(ViennaPSCallableWrapper ${VIENNAPS_GPU_MODELS}/CallableWrapper.cu)
14+
viennacore_add_optixir(
15+
ViennaPSCallableWrapper
16+
${VIENNAPS_GPU_MODELS}/CallableWrapper.cu
17+
DEPENDS
18+
"${VIENNAPS_GPU_MODELS}/FaradayCage.cuh"
19+
"${VIENNAPS_GPU_MODELS}/IonBeamEtching.cuh"
20+
"${VIENNAPS_GPU_MODELS}/MultiParticle.cuh"
21+
"${VIENNAPS_GPU_MODELS}/PlasmaEtching.cuh"
22+
"${VIENNAPS_GPU_MODELS}/SingleParticle.cuh"
23+
"${VIENNAPS_GPU_MODELS}/SingleParticleALD.cuh"
24+
"${VIENNAPS_GPU_MODELS}/TEOSPECVD.cuh")
1525

1626
set(VIENNAPS_GPU_DEPENDENCIES
1727
ViennaPSCallableWrapper ${VIENNARAY_GPU_DEPENDENCIES}

‎gpu/models/PlasmaEtching.cuh‎

Lines changed: 13 additions & 12 deletions
Original file line numberDiff line numberDiff line change
@@ -33,8 +33,8 @@ plasmaNeutralReflection(const void *sbtData, viennaray::gpu::PerRayData *prd) {
3333
float *coverages = (float *)baseData->cellData;
3434
const auto &phi_E = coverages[prd->primID];
3535
const auto &phi_P = coverages[prd->primID + launchParams.numElements];
36-
int material = launchParams.materialIds[prd->primID];
37-
float sticking = launchParams.materialSticking[material];
36+
int id = launchParams.materialIds[prd->primID]; // consecutive ID, not enum
37+
float sticking = launchParams.materialSticking[id];
3838
float Seff = sticking * max(1.f - phi_E - phi_P, 0.f);
3939
prd->rayWeight -= prd->rayWeight * Seff;
4040
auto geoNormal = viennaray::gpu::getNormal(sbtData, prd->primID);
@@ -63,11 +63,12 @@ plasmaNeutralReflectionNoPassivation(const void *sbtData,
6363

6464
__forceinline__ __device__ void
6565
plasmaIonCollision(const void *sbtData, viennaray::gpu::PerRayData *prd) {
66-
viennaps::PlasmaEtchingParameters<float> *params =
67-
reinterpret_cast<viennaps::PlasmaEtchingParameters<float> *>(
66+
viennaps::PlasmaEtchingParametersGPU *params =
67+
reinterpret_cast<viennaps::PlasmaEtchingParametersGPU *>(
6868
launchParams.customData);
6969
for (int i = 0; i < prd->ISCount; ++i) {
70-
int material = launchParams.materialIds[prd->primIDs[i]];
70+
int id = launchParams.materialIds[prd->primIDs[i]]; // consecutive ID
71+
int material = launchParams.materialMap[id]; // mapped to enum
7172
auto geomNormal = viennaray::gpu::getNormal(sbtData, prd->primIDs[i]);
7273
auto cosTheta = __saturatef(
7374
-viennacore::DotProduct(prd->dir, geomNormal)); // clamp to [0,1]
@@ -104,22 +105,22 @@ plasmaIonCollision(const void *sbtData, viennaray::gpu::PerRayData *prd) {
104105
atomicAdd(&launchParams
105106
.resultBuffer[viennaray::gpu::getIdxOffset(0, launchParams) +
106107
prd->primIDs[i]],
107-
Y_sp * prd->rayWeight);
108+
static_cast<viennaray::gpu::ResultType>(Y_sp * prd->rayWeight));
108109
atomicAdd(&launchParams
109110
.resultBuffer[viennaray::gpu::getIdxOffset(1, launchParams) +
110111
prd->primIDs[i]],
111-
Y_Si * prd->rayWeight);
112+
static_cast<viennaray::gpu::ResultType>(Y_Si * prd->rayWeight));
112113
atomicAdd(&launchParams
113114
.resultBuffer[viennaray::gpu::getIdxOffset(2, launchParams) +
114115
prd->primIDs[i]],
115-
Y_P * prd->rayWeight);
116+
static_cast<viennaray::gpu::ResultType>(Y_P * prd->rayWeight));
116117
}
117118
}
118119

119120
__forceinline__ __device__ void
120121
plasmaIonReflection(const void *sbtData, viennaray::gpu::PerRayData *prd) {
121-
viennaps::PlasmaEtchingParameters<float> *params =
122-
reinterpret_cast<viennaps::PlasmaEtchingParameters<float> *>(
122+
viennaps::PlasmaEtchingParametersGPU *params =
123+
reinterpret_cast<viennaps::PlasmaEtchingParametersGPU *>(
123124
launchParams.customData);
124125
auto geomNormal = viennaray::gpu::getNormal(sbtData, prd->primID);
125126
auto cosTheta = __saturatef(
@@ -153,8 +154,8 @@ plasmaIonReflection(const void *sbtData, viennaray::gpu::PerRayData *prd) {
153154
}
154155

155156
__forceinline__ __device__ void plasmaIonInit(viennaray::gpu::PerRayData *prd) {
156-
viennaps::PlasmaEtchingParameters<float> *params =
157-
reinterpret_cast<viennaps::PlasmaEtchingParameters<float> *>(
157+
viennaps::PlasmaEtchingParametersGPU *params =
158+
reinterpret_cast<viennaps::PlasmaEtchingParametersGPU *>(
158159
launchParams.customData);
159160
viennaps::gpu::impl::initNormalDistEnergy(prd, params->Ions.meanEnergy,
160161
params->Ions.sigmaEnergy);

‎include/viennaps/models/psHBrO2Etching.hpp‎

Lines changed: 21 additions & 7 deletions
Original file line numberDiff line numberDiff line change
@@ -22,7 +22,7 @@ template <typename NumericType, int D>
2222
class HBrO2Etching final : public ProcessModelGPU<NumericType, D> {
2323
public:
2424
explicit HBrO2Etching(const PlasmaEtchingParameters<NumericType> &pParams)
25-
: params(pParams), deviceParams(pParams.convertToFloat()) {
25+
: params(pParams), deviceParams(pParams) {
2626
initializeModel();
2727
}
2828

@@ -87,22 +87,34 @@ class HBrO2Etching final : public ProcessModelGPU<NumericType, D> {
8787
"__direct_callable__plasmaNeutralReflection"}};
8888
this->setParticleCallableMap(pMap, cMap);
8989

90-
this->processData.alloc(sizeof(PlasmaEtchingParameters<float>));
90+
this->setUseMaterialIds(true);
91+
precomputeSqrtEnergies();
92+
this->processData.alloc(sizeof(PlasmaEtchingParametersGPU));
9193
this->processData.upload(&deviceParams, 1);
94+
this->hasGPU = true;
9295

93-
this->setUseMaterialIds(true);
9496
this->processMetaData = params.toProcessMetaData();
9597
}
9698

9799
void setParameters(const PlasmaEtchingParameters<NumericType> &pParams) {
98100
params = pParams;
99-
deviceParams = pParams.convertToFloat();
101+
deviceParams = PlasmaEtchingParametersGPU(pParams);
102+
precomputeSqrtEnergies();
100103
this->processData.upload(&deviceParams, 1);
101104
}
102105

103106
private:
104107
PlasmaEtchingParameters<NumericType> params;
105-
PlasmaEtchingParameters<float> deviceParams;
108+
PlasmaEtchingParametersGPU deviceParams;
109+
110+
void precomputeSqrtEnergies() {
111+
deviceParams.Substrate.Eth_ie = std::sqrt(deviceParams.Substrate.Eth_ie);
112+
deviceParams.Passivation.Eth_ie =
113+
std::sqrt(deviceParams.Passivation.Eth_ie);
114+
deviceParams.Substrate.Eth_sp = std::sqrt(deviceParams.Substrate.Eth_sp);
115+
deviceParams.Mask.Eth_sp = std::sqrt(deviceParams.Mask.Eth_sp);
116+
deviceParams.Polymer.Eth_sp = std::sqrt(deviceParams.Polymer.Eth_sp);
117+
}
106118
};
107119
} // namespace gpu
108120
#endif
@@ -164,8 +176,10 @@ class HBrO2Etching : public ProcessModelCPU<NumericType, D> {
164176
defParams.passivationFlux = 1.0e2;
165177

166178
// sticking probabilities
167-
defParams.beta_E = {{1, 0.1}, {0, 0.1}};
168-
defParams.beta_P = {{1, 1.}, {0, 1.}};
179+
defParams.beta_E = {{static_cast<int>(Material::Si), 0.1},
180+
{static_cast<int>(Material::Mask), 0.1}};
181+
defParams.beta_P = {{static_cast<int>(Material::Si), 1.},
182+
{static_cast<int>(Material::Mask), 1.}};
169183

170184
defParams.etchStopDepth = std::numeric_limits<NumericType>::lowest();
171185

‎include/viennaps/models/psPlasmaEtchingParameters.hpp‎

Lines changed: 45 additions & 55 deletions
Original file line numberDiff line numberDiff line change
@@ -4,8 +4,6 @@
44

55
#include <unordered_map>
66

7-
#define FLOAT_CAST(dest, x) dest.x = static_cast<float>(x);
8-
97
namespace viennaps {
108

119
template <typename NumericType> struct PlasmaEtchingParameters {
@@ -148,61 +146,53 @@ template <typename NumericType> struct PlasmaEtchingParameters {
148146

149147
return processData;
150148
}
149+
};
151150

152-
PlasmaEtchingParameters<float> convertToFloat() const {
153-
PlasmaEtchingParameters<float> pParams;
154-
if constexpr (std::is_same_v<NumericType, float>) {
155-
pParams = *this;
156-
} else {
157-
FLOAT_CAST(pParams, ionFlux);
158-
FLOAT_CAST(pParams, etchantFlux);
159-
FLOAT_CAST(pParams, passivationFlux);
160-
161-
for (auto &pair : beta_E) {
162-
pParams.beta_E[pair.first] = static_cast<float>(pair.second);
163-
}
164-
for (auto &pair : beta_P) {
165-
pParams.beta_P[pair.first] = static_cast<float>(pair.second);
166-
}
167-
168-
FLOAT_CAST(pParams, etchStopDepth);
169-
170-
FLOAT_CAST(pParams, Mask.A_sp);
171-
FLOAT_CAST(pParams, Mask.B_sp);
172-
FLOAT_CAST(pParams, Mask.Eth_sp);
173-
FLOAT_CAST(pParams, Mask.rho);
174-
175-
FLOAT_CAST(pParams, Polymer.A_sp);
176-
FLOAT_CAST(pParams, Polymer.B_sp);
177-
FLOAT_CAST(pParams, Polymer.Eth_sp);
178-
FLOAT_CAST(pParams, Polymer.rho);
179-
180-
FLOAT_CAST(pParams, Substrate.A_ie);
181-
FLOAT_CAST(pParams, Substrate.A_sp);
182-
FLOAT_CAST(pParams, Substrate.B_ie);
183-
FLOAT_CAST(pParams, Substrate.B_sp);
184-
FLOAT_CAST(pParams, Substrate.Eth_ie);
185-
FLOAT_CAST(pParams, Substrate.Eth_sp);
186-
FLOAT_CAST(pParams, Substrate.k_sigma);
187-
FLOAT_CAST(pParams, Substrate.beta_sigma);
188-
FLOAT_CAST(pParams, Substrate.rho);
189-
190-
FLOAT_CAST(pParams, Passivation.A_ie);
191-
FLOAT_CAST(pParams, Passivation.Eth_ie);
192-
193-
FLOAT_CAST(pParams, Ions.exponent);
194-
FLOAT_CAST(pParams, Ions.meanEnergy);
195-
FLOAT_CAST(pParams, Ions.sigmaEnergy);
196-
FLOAT_CAST(pParams, Ions.inflectAngle);
197-
FLOAT_CAST(pParams, Ions.n_l);
198-
FLOAT_CAST(pParams, Ions.minAngle);
199-
FLOAT_CAST(pParams, Ions.thetaRMin);
200-
FLOAT_CAST(pParams, Ions.thetaRMax);
201-
}
202-
return pParams;
151+
#ifdef VIENNACORE_COMPILE_GPU
152+
struct PlasmaEtchingParametersGPU {
153+
PlasmaEtchingParametersGPU() = default;
154+
template <typename NumericType>
155+
PlasmaEtchingParametersGPU(
156+
const PlasmaEtchingParameters<NumericType> &parameters) {
157+
Mask.A_sp = static_cast<float>(parameters.Mask.A_sp);
158+
Mask.B_sp = static_cast<float>(parameters.Mask.B_sp);
159+
Mask.Eth_sp = static_cast<float>(parameters.Mask.Eth_sp);
160+
Mask.rho = static_cast<float>(parameters.Mask.rho);
161+
162+
Polymer.A_sp = static_cast<float>(parameters.Polymer.A_sp);
163+
Polymer.B_sp = static_cast<float>(parameters.Polymer.B_sp);
164+
Polymer.Eth_sp = static_cast<float>(parameters.Polymer.Eth_sp);
165+
Polymer.rho = static_cast<float>(parameters.Polymer.rho);
166+
167+
Substrate.rho = static_cast<float>(parameters.Substrate.rho);
168+
Substrate.Eth_sp = static_cast<float>(parameters.Substrate.Eth_sp);
169+
Substrate.Eth_ie = static_cast<float>(parameters.Substrate.Eth_ie);
170+
Substrate.A_sp = static_cast<float>(parameters.Substrate.A_sp);
171+
Substrate.B_sp = static_cast<float>(parameters.Substrate.B_sp);
172+
Substrate.A_ie = static_cast<float>(parameters.Substrate.A_ie);
173+
Substrate.B_ie = static_cast<float>(parameters.Substrate.B_ie);
174+
Substrate.k_sigma = static_cast<float>(parameters.Substrate.k_sigma);
175+
Substrate.beta_sigma = static_cast<float>(parameters.Substrate.beta_sigma);
176+
177+
Passivation.Eth_ie = static_cast<float>(parameters.Passivation.Eth_ie);
178+
Passivation.A_ie = static_cast<float>(parameters.Passivation.A_ie);
179+
180+
Ions.meanEnergy = static_cast<float>(parameters.Ions.meanEnergy);
181+
Ions.sigmaEnergy = static_cast<float>(parameters.Ions.sigmaEnergy);
182+
Ions.exponent = static_cast<float>(parameters.Ions.exponent);
183+
Ions.inflectAngle = static_cast<float>(parameters.Ions.inflectAngle);
184+
Ions.n_l = static_cast<float>(parameters.Ions.n_l);
185+
Ions.minAngle = static_cast<float>(parameters.Ions.minAngle);
186+
Ions.thetaRMin = static_cast<float>(parameters.Ions.thetaRMin);
187+
Ions.thetaRMax = static_cast<float>(parameters.Ions.thetaRMax);
203188
}
189+
190+
PlasmaEtchingParameters<float>::MaskType Mask;
191+
PlasmaEtchingParameters<float>::PolymerType Polymer;
192+
PlasmaEtchingParameters<float>::MaterialType Substrate;
193+
PlasmaEtchingParameters<float>::PassivationType Passivation;
194+
PlasmaEtchingParameters<float>::IonType Ions;
204195
};
196+
#endif
205197

206198
} // namespace viennaps
207-
208-
#undef FLOAT_CAST

0 commit comments

Comments
 (0)