Add GPU support for SF6-C4F8 plasma model

gpu/include/models/CallableWrapper.cu

@@ -93,6 +93,12 @@ __direct_callable__plasmaNeutralReflection(const void *sbtData,
   plasmaNeutralReflection(sbtData, prd);
 }
 
+extern "C" __device__ void
+__direct_callable__plasmaNeutralReflectionNoPassivation(
+    const void *sbtData, viennaray::gpu::PerRayData *prd) {
+  plasmaNeutralReflectionNoPassivation(sbtData, prd);
+}
+
 extern "C" __device__ void
 __direct_callable__plasmaIonCollision(const void *sbtData,
                                       viennaray::gpu::PerRayData *prd) {

gpu/include/models/PlasmaEtching.cuh

@@ -41,6 +41,22 @@ plasmaNeutralReflection(const void *sbtData, viennaray::gpu::PerRayData *prd) {
   viennaray::gpu::diffuseReflection(prd, geoNormal, launchParams.D);
 }
 
+// Specialized neutral reflection for models without passivation (e.g., SF6C4F8)
+__forceinline__ __device__ void
+plasmaNeutralReflectionNoPassivation(const void *sbtData,
+                                      viennaray::gpu::PerRayData *prd) {
+  const viennaray::gpu::HitSBTDataBase *baseData =
+      reinterpret_cast<const viennaray::gpu::HitSBTDataBase *>(sbtData);
+  float *data = (float *)baseData->cellData;
+  const auto &phi_E = data[prd->primID];
+  int material = launchParams.materialIds[prd->primID];
+  float sticking = launchParams.materialSticking[material];
+  float Seff = sticking * max(1.f - phi_E, 0.f);
+  prd->rayWeight -= prd->rayWeight * Seff;
+  auto geoNormal = viennaray::gpu::getNormal(sbtData, prd->primID);
+  viennaray::gpu::diffuseReflection(prd, geoNormal, launchParams.D);
+}
+
 //
 // --- Ion particle
 //

gpu/include/psgPipelineParameters.hpp

@@ -115,4 +115,4 @@ initNormalDistEnergy(viennaray::gpu::PerRayData *prd, const float mean,
 }
 #endif
 
-} // namespace viennaps::gpu::impl
+} // namespace viennaps::gpu::impl

include/viennaps/models/psSF6C4F8Etching.hpp

@@ -17,6 +17,101 @@ namespace viennaps {
 
 using namespace viennacore;
 
+#ifdef VIENNACORE_COMPILE_GPU
+namespace gpu {
+/// GPU Version of the SF6/C4F8 plasma etching model
+template <typename NumericType, int D>
+class SF6C4F8Etching final : public ProcessModelGPU<NumericType, D> {
+public:
+  explicit SF6C4F8Etching(const PlasmaEtchingParameters<NumericType> &pParams)
+      : params(pParams), deviceParams(pParams.convertToFloat()) {
+    initializeModel();
+  }
+
+  ~SF6C4F8Etching() override { this->processData.free(); }
+
+private:
+  void initializeModel() {
+    // particles
+    viennaray::gpu::Particle<NumericType> ion;
+    ion.name = "Ion"; // name for shader programs postfix
+    ion.dataLabels.push_back("ionSputterFlux");
+    ion.dataLabels.push_back("ionEnhancedFlux");
+    ion.dataLabels.push_back("ionEnhancedPassivationFlux");
+    ion.sticking = 0.f;
+    ion.cosineExponent = params.Ions.exponent;
+
+    viennaray::gpu::Particle<NumericType> etchant;
+    etchant.name = "Etchant";
+    etchant.dataLabels.push_back("etchantFlux");
+    etchant.cosineExponent = 1.f;
+    etchant.materialSticking = params.beta_E;
+
+    // No oxygen/passivation particle for SF6C4F8
+
+    // surface model
+    auto surfModel = SmartPointer<
+        viennaps::impl::PlasmaEtchingSurfaceModel<NumericType, D>>::New(params);
+
+    // velocity field
+    auto velField = SmartPointer<DefaultVelocityField<NumericType, D>>::New();
+
+    this->setSurfaceModel(surfModel);
+    this->setVelocityField(velField);
+    this->setProcessName("SF6C4F8Etching");
+    this->getParticleTypes().clear();
+
+    this->insertNextParticleType(ion);
+    this->insertNextParticleType(etchant);
+
+    std::unordered_map<std::string, unsigned> pMap = {{"Ion", 0},
+                                                      {"Etchant", 1}};
+    std::vector<viennaray::gpu::CallableConfig> cMap = {
+        {0, viennaray::gpu::CallableSlot::COLLISION,
+         "__direct_callable__plasmaIonCollision"},
+        {0, viennaray::gpu::CallableSlot::REFLECTION,
+         "__direct_callable__plasmaIonReflection"},
+        {0, viennaray::gpu::CallableSlot::INIT,
+         "__direct_callable__plasmaIonInit"},
+        {1, viennaray::gpu::CallableSlot::COLLISION,
+         "__direct_callable__plasmaNeutralCollision"},
+        {1, viennaray::gpu::CallableSlot::REFLECTION,
+         "__direct_callable__plasmaNeutralReflectionNoPassivation"}};
+    this->setParticleCallableMap(pMap, cMap);
+    this->setCallableFileName("CallableWrapper");
+
+    this->setUseMaterialIds(true);
+    precomputeSqrtEnergies();
+    this->processData.alloc(sizeof(PlasmaEtchingParameters<float>));
+    this->processData.upload(&deviceParams, 1);
+    this->hasGPU = true;
+
+    this->processMetaData = params.toProcessMetaData();
+  }
+
+  void setParameters(const PlasmaEtchingParameters<NumericType> &pParams) {
+    params = pParams;
+    deviceParams = pParams.convertToFloat();
+    precomputeSqrtEnergies();
+    this->processData.upload(&deviceParams, 1);
+  }
+
+private:
+  PlasmaEtchingParameters<NumericType> params;
+  PlasmaEtchingParameters<float> deviceParams;
+
+  void precomputeSqrtEnergies() {
+    deviceParams.Substrate.Eth_ie = std::sqrt(deviceParams.Substrate.Eth_ie);
+    deviceParams.Passivation.Eth_ie =
+        std::sqrt(deviceParams.Passivation.Eth_ie);
+    deviceParams.Substrate.Eth_sp = std::sqrt(deviceParams.Substrate.Eth_sp);
+    deviceParams.Mask.Eth_sp = std::sqrt(deviceParams.Mask.Eth_sp);
+    deviceParams.Polymer.Eth_sp = std::sqrt(deviceParams.Polymer.Eth_sp);
+  }
+};
+} // namespace gpu
+#endif
+
 /// Model for etching Si in a SF6/C4F8 plasma. This model extends the SF6O2
 /// etching model to include polymer etching. The polymer layer is deposited
 /// by a separate process and etched by sputtering similarly to the mask
@@ -50,6 +145,12 @@ class SF6C4F8Etching : public ProcessModelCPU<NumericType, D> {
     initializeModel();
   }
 
+#ifdef VIENNACORE_COMPILE_GPU
+  SmartPointer<ProcessModelBase<NumericType, D>> getGPUModel() override {
+    return SmartPointer<gpu::SF6C4F8Etching<NumericType, D>>::New(params);
+  }
+#endif
+
   void setParameters(const PlasmaEtchingParameters<NumericType> &pParams) {
     params = pParams;
     initializeModel();
@@ -147,6 +248,7 @@ class SF6C4F8Etching : public ProcessModelCPU<NumericType, D> {
     this->setVelocityField(velField);
 
     this->setProcessName("SF6C4F8Etching");
+    this->hasGPU = true;
 
     processMetaData = params.toProcessMetaData();
     // add units

include/viennaps/process/psFluxProcessStrategy.hpp

@@ -157,18 +157,19 @@ class FluxProcessStrategy : public ProcessStrategy<NumericType, D> {
 
     // Try to initialize coverages
     meshGenerator_.apply();
-    if (!context.coverageParams.initialized) {
-      if (coverageManager_.initializeCoverages(context)) {
+    // Always reinitialize coverages to handle geometry changes between cycles
+    if (coverageManager_.initializeCoverages(context)) {
+      if (!context.coverageParams.initialized) {
         context.flags.useCoverages = true;
         Logger::getInstance().addInfo("Using coverages.").print();
 
         // Translator is needed to map coverage values to level set points
         if (!translator_)
           translator_ = SmartPointer<TranslatorType>::New();
         meshGenerator_.setTranslator(translator_);
+      } else {
+        Logger::getInstance().addInfo("Coverages reinitialized.").print();
       }
-    } else {
-      Logger::getInstance().addInfo("Coverages already initialized.").print();
     }
     context.model->getSurfaceModel()->initializeSurfaceData(
         context.diskMesh->nodes.size());