Skip to content

Commit 9139226

Browse files
kenyastylekenyastyle
authored
Add GPU support for SF6-C4F8 plasma model (#179)
* add GPU support for SF6C4F8 etching * add a dedicated reflection model in case of no passivation * Always reinitialize coverages to handle geometry changes between cycles * fix formatting
1 parent c7dcf89 commit 9139226

5 files changed

Lines changed: 130 additions & 5 deletions

File tree

‎gpu/include/models/CallableWrapper.cu‎

Lines changed: 6 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -93,6 +93,12 @@ __direct_callable__plasmaNeutralReflection(const void *sbtData,
9393
plasmaNeutralReflection(sbtData, prd);
9494
}
9595

96+
extern "C" __device__ void
97+
__direct_callable__plasmaNeutralReflectionNoPassivation(
98+
const void *sbtData, viennaray::gpu::PerRayData *prd) {
99+
plasmaNeutralReflectionNoPassivation(sbtData, prd);
100+
}
101+
96102
extern "C" __device__ void
97103
__direct_callable__plasmaIonCollision(const void *sbtData,
98104
viennaray::gpu::PerRayData *prd) {

‎gpu/include/models/PlasmaEtching.cuh‎

Lines changed: 16 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -41,6 +41,22 @@ plasmaNeutralReflection(const void *sbtData, viennaray::gpu::PerRayData *prd) {
4141
viennaray::gpu::diffuseReflection(prd, geoNormal, launchParams.D);
4242
}
4343

44+
// Specialized neutral reflection for models without passivation (e.g., SF6C4F8)
45+
__forceinline__ __device__ void
46+
plasmaNeutralReflectionNoPassivation(const void *sbtData,
47+
viennaray::gpu::PerRayData *prd) {
48+
const viennaray::gpu::HitSBTDataBase *baseData =
49+
reinterpret_cast<const viennaray::gpu::HitSBTDataBase *>(sbtData);
50+
float *data = (float *)baseData->cellData;
51+
const auto &phi_E = data[prd->primID];
52+
int material = launchParams.materialIds[prd->primID];
53+
float sticking = launchParams.materialSticking[material];
54+
float Seff = sticking * max(1.f - phi_E, 0.f);
55+
prd->rayWeight -= prd->rayWeight * Seff;
56+
auto geoNormal = viennaray::gpu::getNormal(sbtData, prd->primID);
57+
viennaray::gpu::diffuseReflection(prd, geoNormal, launchParams.D);
58+
}
59+
4460
//
4561
// --- Ion particle
4662
//

‎gpu/include/psgPipelineParameters.hpp‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -115,4 +115,4 @@ initNormalDistEnergy(viennaray::gpu::PerRayData *prd, const float mean,
115115
}
116116
#endif
117117

118-
} // namespace viennaps::gpu::impl
118+
} // namespace viennaps::gpu::impl

‎include/viennaps/models/psSF6C4F8Etching.hpp‎

Lines changed: 102 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -17,6 +17,101 @@ namespace viennaps {
1717

1818
using namespace viennacore;
1919

20+
#ifdef VIENNACORE_COMPILE_GPU
21+
namespace gpu {
22+
/// GPU Version of the SF6/C4F8 plasma etching model
23+
template <typename NumericType, int D>
24+
class SF6C4F8Etching final : public ProcessModelGPU<NumericType, D> {
25+
public:
26+
explicit SF6C4F8Etching(const PlasmaEtchingParameters<NumericType> &pParams)
27+
: params(pParams), deviceParams(pParams.convertToFloat()) {
28+
initializeModel();
29+
}
30+
31+
~SF6C4F8Etching() override { this->processData.free(); }
32+
33+
private:
34+
void initializeModel() {
35+
// particles
36+
viennaray::gpu::Particle<NumericType> ion;
37+
ion.name = "Ion"; // name for shader programs postfix
38+
ion.dataLabels.push_back("ionSputterFlux");
39+
ion.dataLabels.push_back("ionEnhancedFlux");
40+
ion.dataLabels.push_back("ionEnhancedPassivationFlux");
41+
ion.sticking = 0.f;
42+
ion.cosineExponent = params.Ions.exponent;
43+
44+
viennaray::gpu::Particle<NumericType> etchant;
45+
etchant.name = "Etchant";
46+
etchant.dataLabels.push_back("etchantFlux");
47+
etchant.cosineExponent = 1.f;
48+
etchant.materialSticking = params.beta_E;
49+
50+
// No oxygen/passivation particle for SF6C4F8
51+
52+
// surface model
53+
auto surfModel = SmartPointer<
54+
viennaps::impl::PlasmaEtchingSurfaceModel<NumericType, D>>::New(params);
55+
56+
// velocity field
57+
auto velField = SmartPointer<DefaultVelocityField<NumericType, D>>::New();
58+
59+
this->setSurfaceModel(surfModel);
60+
this->setVelocityField(velField);
61+
this->setProcessName("SF6C4F8Etching");
62+
this->getParticleTypes().clear();
63+
64+
this->insertNextParticleType(ion);
65+
this->insertNextParticleType(etchant);
66+
67+
std::unordered_map<std::string, unsigned> pMap = {{"Ion", 0},
68+
{"Etchant", 1}};
69+
std::vector<viennaray::gpu::CallableConfig> cMap = {
70+
{0, viennaray::gpu::CallableSlot::COLLISION,
71+
"__direct_callable__plasmaIonCollision"},
72+
{0, viennaray::gpu::CallableSlot::REFLECTION,
73+
"__direct_callable__plasmaIonReflection"},
74+
{0, viennaray::gpu::CallableSlot::INIT,
75+
"__direct_callable__plasmaIonInit"},
76+
{1, viennaray::gpu::CallableSlot::COLLISION,
77+
"__direct_callable__plasmaNeutralCollision"},
78+
{1, viennaray::gpu::CallableSlot::REFLECTION,
79+
"__direct_callable__plasmaNeutralReflectionNoPassivation"}};
80+
this->setParticleCallableMap(pMap, cMap);
81+
this->setCallableFileName("CallableWrapper");
82+
83+
this->setUseMaterialIds(true);
84+
precomputeSqrtEnergies();
85+
this->processData.alloc(sizeof(PlasmaEtchingParameters<float>));
86+
this->processData.upload(&deviceParams, 1);
87+
this->hasGPU = true;
88+
89+
this->processMetaData = params.toProcessMetaData();
90+
}
91+
92+
void setParameters(const PlasmaEtchingParameters<NumericType> &pParams) {
93+
params = pParams;
94+
deviceParams = pParams.convertToFloat();
95+
precomputeSqrtEnergies();
96+
this->processData.upload(&deviceParams, 1);
97+
}
98+
99+
private:
100+
PlasmaEtchingParameters<NumericType> params;
101+
PlasmaEtchingParameters<float> deviceParams;
102+
103+
void precomputeSqrtEnergies() {
104+
deviceParams.Substrate.Eth_ie = std::sqrt(deviceParams.Substrate.Eth_ie);
105+
deviceParams.Passivation.Eth_ie =
106+
std::sqrt(deviceParams.Passivation.Eth_ie);
107+
deviceParams.Substrate.Eth_sp = std::sqrt(deviceParams.Substrate.Eth_sp);
108+
deviceParams.Mask.Eth_sp = std::sqrt(deviceParams.Mask.Eth_sp);
109+
deviceParams.Polymer.Eth_sp = std::sqrt(deviceParams.Polymer.Eth_sp);
110+
}
111+
};
112+
} // namespace gpu
113+
#endif
114+
20115
/// Model for etching Si in a SF6/C4F8 plasma. This model extends the SF6O2
21116
/// etching model to include polymer etching. The polymer layer is deposited
22117
/// by a separate process and etched by sputtering similarly to the mask
@@ -50,6 +145,12 @@ class SF6C4F8Etching : public ProcessModelCPU<NumericType, D> {
50145
initializeModel();
51146
}
52147

148+
#ifdef VIENNACORE_COMPILE_GPU
149+
SmartPointer<ProcessModelBase<NumericType, D>> getGPUModel() override {
150+
return SmartPointer<gpu::SF6C4F8Etching<NumericType, D>>::New(params);
151+
}
152+
#endif
153+
53154
void setParameters(const PlasmaEtchingParameters<NumericType> &pParams) {
54155
params = pParams;
55156
initializeModel();
@@ -147,6 +248,7 @@ class SF6C4F8Etching : public ProcessModelCPU<NumericType, D> {
147248
this->setVelocityField(velField);
148249

149250
this->setProcessName("SF6C4F8Etching");
251+
this->hasGPU = true;
150252

151253
processMetaData = params.toProcessMetaData();
152254
// add units

‎include/viennaps/process/psFluxProcessStrategy.hpp‎

Lines changed: 5 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -157,18 +157,19 @@ class FluxProcessStrategy : public ProcessStrategy<NumericType, D> {
157157

158158
// Try to initialize coverages
159159
meshGenerator_.apply();
160-
if (!context.coverageParams.initialized) {
161-
if (coverageManager_.initializeCoverages(context)) {
160+
// Always reinitialize coverages to handle geometry changes between cycles
161+
if (coverageManager_.initializeCoverages(context)) {
162+
if (!context.coverageParams.initialized) {
162163
context.flags.useCoverages = true;
163164
Logger::getInstance().addInfo("Using coverages.").print();
164165

165166
// Translator is needed to map coverage values to level set points
166167
if (!translator_)
167168
translator_ = SmartPointer<TranslatorType>::New();
168169
meshGenerator_.setTranslator(translator_);
170+
} else {
171+
Logger::getInstance().addInfo("Coverages reinitialized.").print();
169172
}
170-
} else {
171-
Logger::getInstance().addInfo("Coverages already initialized.").print();
172173
}
173174
context.model->getSurfaceModel()->initializeSurfaceData(
174175
context.diskMesh->nodes.size());

0 commit comments

Comments
 (0)