fix: HBrO2 model configuration and material handling on GPU

CMakeLists.txt

@@ -105,7 +105,7 @@ include("cmake/cpm.cmake")
 
 CPMAddPackage(
   NAME ViennaCore
-  VERSION 1.9.1
+  VERSION 1.9.2
   GIT_REPOSITORY "https://github.com/ViennaTools/ViennaCore"
   EXCLUDE_FROM_ALL ${VIENNAPS_BUILD_PYTHON}
   OPTIONS "VIENNACORE_USE_GPU ${VIENNAPS_USE_GPU}")
@@ -117,7 +117,7 @@ CPMAddPackage(
 
 CPMAddPackage(
   NAME ViennaRay
-  VERSION 3.10.0
+  VERSION 3.10.1
   GIT_REPOSITORY "https://github.com/ViennaTools/ViennaRay"
   EXCLUDE_FROM_ALL ${VIENNAPS_BUILD_PYTHON}
   OPTIONS "VIENNARAY_USE_GPU ${VIENNAPS_USE_GPU}")

examples/DRAMWiggling/DRAMWiggling.cpp

@@ -6,9 +6,6 @@ int main(int argc, char **argv) {
   using NumericType = double;
   constexpr int D = 3;
 
-  Logger::setLogLevel(LogLevel::ERROR);
-  omp_set_num_threads(12);
-
   // Parse the parameters
   util::Parameters params;
   if (argc > 1) {
@@ -51,6 +48,7 @@ int main(int argc, char **argv) {
   modelParams.Ions.sigmaEnergy = params.get("sigmaEnergy");
   modelParams.Ions.exponent = params.get("ionExponent");
   modelParams.Ions.n_l = 200;
+  modelParams.Substrate.B_sp = 0.75;
   auto model = SmartPointer<HBrO2Etching<NumericType, D>>::New(modelParams);
 
   // Advection parameters

examples/DRAMWiggling/DRAMWiggling.py

@@ -1,4 +1,3 @@
-import viennaps.d3 as psd
 import viennaps as ps
 from argparse import ArgumentParser
 
@@ -10,6 +9,8 @@
 parser.add_argument("filename")
 args = parser.parse_args()
 
+ps.setDimension(3)
+
 gridDelta = 0.01 * (1.0 + 1e-12)
 boundaryConds = [
     ps.BoundaryType.REFLECTIVE_BOUNDARY,
@@ -19,36 +20,34 @@
 
 params = ps.readConfigFile(args.filename)
 
-mask = psd.GDSGeometry(gridDelta, boundaryConds)
+mask = ps.GDSGeometry(gridDelta, boundaryConds)
 mask.setBoundaryPadding(0.1, 0.1)
-reader = psd.GDSReader(mask, params["gdsFile"])
+reader = ps.GDSReader(mask, params["gdsFile"])
 reader.apply()
 
 # Prepare geometry
-geometry = psd.Domain()
+geometry = ps.Domain()
 
 # Insert GDS layers
 maskLS = mask.layerToLevelSet(0, 0.0, 0.18)
 geometry.insertNextLevelSetAsMaterial(maskLS, ps.Material.Mask)
 
 # Add plane
-psd.MakePlane(geometry, 0.0, ps.Material.Si, True).apply()
-
-# print intermediate output surfaces during the process
-ps.Logger.setLogLevel(ps.LogLevel.INFO)
+ps.MakePlane(geometry, 0.0, ps.Material.Si, True).apply()
 
 ps.Length.setUnit(params["lengthUnit"])
 ps.Time.setUnit(params["timeUnit"])
 
-modelParams = psd.HBrO2Etching.defaultParameters()
+modelParams = ps.HBrO2Etching.defaultParameters()
 modelParams.ionFlux = params["ionFlux"]
 modelParams.etchantFlux = params["etchantFlux"]
 modelParams.passivationFlux = params["oxygenFlux"]
 modelParams.Ions.meanEnergy = params["meanEnergy"]
 modelParams.Ions.sigmaEnergy = params["sigmaEnergy"]
 modelParams.Ions.exponent = params["ionExponent"]
 modelParams.Ions.n_l = 200
-model = psd.HBrO2Etching(modelParams)
+modelParams.Substrate.B_sp = 0.75
+model = ps.HBrO2Etching(modelParams)
 
 coverageParameters = ps.CoverageParameters()
 coverageParameters.maxIterations = 10
@@ -57,12 +56,10 @@
 rayTracingParams.raysPerPoint = int(params["raysPerPoint"])
 
 advectionParams = ps.AdvectionParameters()
-advectionParams.spatialScheme = ps.util.convertSpatialScheme(
-    params["spatialScheme"]
-)
+advectionParams.spatialScheme = ps.util.convertSpatialScheme(params["spatialScheme"])
 
 # process setup
-process = psd.Process(geometry, model)
+process = ps.Process(geometry, model)
 process.setProcessDuration(params["processTime"])  # seconds
 process.setParameters(coverageParameters)
 process.setParameters(rayTracingParams)

examples/DRAMWiggling/config.txt

@@ -19,6 +19,6 @@ sigmaEnergy=10 # eV
 
 spatialScheme=LF_2
 
-numSteps=100
+numSteps=20
 raysPerPoint=1000
 gdsFile=wiggle_full.gds

gpu/CMakeLists.txt

@@ -11,7 +11,17 @@ if(NOT VIENNAPS_PASSED_FIRST_CONFIGURE)
 endif()
 
 # Add the callable wrapper
-viennacore_add_optixir(ViennaPSCallableWrapper ${VIENNAPS_GPU_MODELS}/CallableWrapper.cu)
+viennacore_add_optixir(
+  ViennaPSCallableWrapper
+  ${VIENNAPS_GPU_MODELS}/CallableWrapper.cu
+  DEPENDS
+  "${VIENNAPS_GPU_MODELS}/FaradayCage.cuh"
+  "${VIENNAPS_GPU_MODELS}/IonBeamEtching.cuh"
+  "${VIENNAPS_GPU_MODELS}/MultiParticle.cuh"
+  "${VIENNAPS_GPU_MODELS}/PlasmaEtching.cuh"
+  "${VIENNAPS_GPU_MODELS}/SingleParticle.cuh"
+  "${VIENNAPS_GPU_MODELS}/SingleParticleALD.cuh"
+  "${VIENNAPS_GPU_MODELS}/TEOSPECVD.cuh")
 
 set(VIENNAPS_GPU_DEPENDENCIES
     ViennaPSCallableWrapper ${VIENNARAY_GPU_DEPENDENCIES}

gpu/models/PlasmaEtching.cuh

@@ -33,8 +33,8 @@ plasmaNeutralReflection(const void *sbtData, viennaray::gpu::PerRayData *prd) {
   float *coverages = (float *)baseData->cellData;
   const auto &phi_E = coverages[prd->primID];
   const auto &phi_P = coverages[prd->primID + launchParams.numElements];
-  int material = launchParams.materialIds[prd->primID];
-  float sticking = launchParams.materialSticking[material];
+  int id = launchParams.materialIds[prd->primID]; // consecutive ID, not enum
+  float sticking = launchParams.materialSticking[id];
   float Seff = sticking * max(1.f - phi_E - phi_P, 0.f);
   prd->rayWeight -= prd->rayWeight * Seff;
   auto geoNormal = viennaray::gpu::getNormal(sbtData, prd->primID);
@@ -63,11 +63,12 @@ plasmaNeutralReflectionNoPassivation(const void *sbtData,
 
 __forceinline__ __device__ void
 plasmaIonCollision(const void *sbtData, viennaray::gpu::PerRayData *prd) {
-  viennaps::PlasmaEtchingParameters<float> *params =
-      reinterpret_cast<viennaps::PlasmaEtchingParameters<float> *>(
+  viennaps::PlasmaEtchingParametersGPU *params =
+      reinterpret_cast<viennaps::PlasmaEtchingParametersGPU *>(
           launchParams.customData);
   for (int i = 0; i < prd->ISCount; ++i) {
-    int material = launchParams.materialIds[prd->primIDs[i]];
+    int id = launchParams.materialIds[prd->primIDs[i]]; // consecutive ID
+    int material = launchParams.materialMap[id];        // mapped to enum
     auto geomNormal = viennaray::gpu::getNormal(sbtData, prd->primIDs[i]);
     auto cosTheta = __saturatef(
         -viennacore::DotProduct(prd->dir, geomNormal)); // clamp to [0,1]
@@ -104,22 +105,22 @@ plasmaIonCollision(const void *sbtData, viennaray::gpu::PerRayData *prd) {
     atomicAdd(&launchParams
                    .resultBuffer[viennaray::gpu::getIdxOffset(0, launchParams) +
                                  prd->primIDs[i]],
-              Y_sp * prd->rayWeight);
+              static_cast<viennaray::gpu::ResultType>(Y_sp * prd->rayWeight));
     atomicAdd(&launchParams
                    .resultBuffer[viennaray::gpu::getIdxOffset(1, launchParams) +
                                  prd->primIDs[i]],
-              Y_Si * prd->rayWeight);
+              static_cast<viennaray::gpu::ResultType>(Y_Si * prd->rayWeight));
     atomicAdd(&launchParams
                    .resultBuffer[viennaray::gpu::getIdxOffset(2, launchParams) +
                                  prd->primIDs[i]],
-              Y_P * prd->rayWeight);
+              static_cast<viennaray::gpu::ResultType>(Y_P * prd->rayWeight));
   }
 }
 
 __forceinline__ __device__ void
 plasmaIonReflection(const void *sbtData, viennaray::gpu::PerRayData *prd) {
-  viennaps::PlasmaEtchingParameters<float> *params =
-      reinterpret_cast<viennaps::PlasmaEtchingParameters<float> *>(
+  viennaps::PlasmaEtchingParametersGPU *params =
+      reinterpret_cast<viennaps::PlasmaEtchingParametersGPU *>(
           launchParams.customData);
   auto geomNormal = viennaray::gpu::getNormal(sbtData, prd->primID);
   auto cosTheta = __saturatef(
@@ -153,8 +154,8 @@ plasmaIonReflection(const void *sbtData, viennaray::gpu::PerRayData *prd) {
 }
 
 __forceinline__ __device__ void plasmaIonInit(viennaray::gpu::PerRayData *prd) {
-  viennaps::PlasmaEtchingParameters<float> *params =
-      reinterpret_cast<viennaps::PlasmaEtchingParameters<float> *>(
+  viennaps::PlasmaEtchingParametersGPU *params =
+      reinterpret_cast<viennaps::PlasmaEtchingParametersGPU *>(
           launchParams.customData);
   viennaps::gpu::impl::initNormalDistEnergy(prd, params->Ions.meanEnergy,
                                             params->Ions.sigmaEnergy);

include/viennaps/models/psHBrO2Etching.hpp

@@ -22,7 +22,7 @@ template <typename NumericType, int D>
 class HBrO2Etching final : public ProcessModelGPU<NumericType, D> {
 public:
   explicit HBrO2Etching(const PlasmaEtchingParameters<NumericType> &pParams)
-      : params(pParams), deviceParams(pParams.convertToFloat()) {
+      : params(pParams), deviceParams(pParams) {
     initializeModel();
   }
 
@@ -87,22 +87,34 @@ class HBrO2Etching final : public ProcessModelGPU<NumericType, D> {
          "__direct_callable__plasmaNeutralReflection"}};
     this->setParticleCallableMap(pMap, cMap);
 
-    this->processData.alloc(sizeof(PlasmaEtchingParameters<float>));
+    this->setUseMaterialIds(true);
+    precomputeSqrtEnergies();
+    this->processData.alloc(sizeof(PlasmaEtchingParametersGPU));
     this->processData.upload(&deviceParams, 1);
+    this->hasGPU = true;
 
-    this->setUseMaterialIds(true);
     this->processMetaData = params.toProcessMetaData();
   }
 
   void setParameters(const PlasmaEtchingParameters<NumericType> &pParams) {
     params = pParams;
-    deviceParams = pParams.convertToFloat();
+    deviceParams = PlasmaEtchingParametersGPU(pParams);
+    precomputeSqrtEnergies();
     this->processData.upload(&deviceParams, 1);
   }
 
 private:
   PlasmaEtchingParameters<NumericType> params;
-  PlasmaEtchingParameters<float> deviceParams;
+  PlasmaEtchingParametersGPU 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
@@ -164,8 +176,10 @@ class HBrO2Etching : public ProcessModelCPU<NumericType, D> {
     defParams.passivationFlux = 1.0e2;
 
     // sticking probabilities
-    defParams.beta_E = {{1, 0.1}, {0, 0.1}};
-    defParams.beta_P = {{1, 1.}, {0, 1.}};
+    defParams.beta_E = {{static_cast<int>(Material::Si), 0.1},
+                        {static_cast<int>(Material::Mask), 0.1}};
+    defParams.beta_P = {{static_cast<int>(Material::Si), 1.},
+                        {static_cast<int>(Material::Mask), 1.}};
 
     defParams.etchStopDepth = std::numeric_limits<NumericType>::lowest();
 

include/viennaps/models/psPlasmaEtchingParameters.hpp

@@ -4,8 +4,6 @@
 
 #include <unordered_map>
 
-#define FLOAT_CAST(dest, x) dest.x = static_cast<float>(x);
-
 namespace viennaps {
 
 template <typename NumericType> struct PlasmaEtchingParameters {
@@ -148,61 +146,53 @@ template <typename NumericType> struct PlasmaEtchingParameters {
 
     return processData;
   }
+};
 
-  PlasmaEtchingParameters<float> convertToFloat() const {
-    PlasmaEtchingParameters<float> pParams;
-    if constexpr (std::is_same_v<NumericType, float>) {
-      pParams = *this;
-    } else {
-      FLOAT_CAST(pParams, ionFlux);
-      FLOAT_CAST(pParams, etchantFlux);
-      FLOAT_CAST(pParams, passivationFlux);
-
-      for (auto &pair : beta_E) {
-        pParams.beta_E[pair.first] = static_cast<float>(pair.second);
-      }
-      for (auto &pair : beta_P) {
-        pParams.beta_P[pair.first] = static_cast<float>(pair.second);
-      }
-
-      FLOAT_CAST(pParams, etchStopDepth);
-
-      FLOAT_CAST(pParams, Mask.A_sp);
-      FLOAT_CAST(pParams, Mask.B_sp);
-      FLOAT_CAST(pParams, Mask.Eth_sp);
-      FLOAT_CAST(pParams, Mask.rho);
-
-      FLOAT_CAST(pParams, Polymer.A_sp);
-      FLOAT_CAST(pParams, Polymer.B_sp);
-      FLOAT_CAST(pParams, Polymer.Eth_sp);
-      FLOAT_CAST(pParams, Polymer.rho);
-
-      FLOAT_CAST(pParams, Substrate.A_ie);
-      FLOAT_CAST(pParams, Substrate.A_sp);
-      FLOAT_CAST(pParams, Substrate.B_ie);
-      FLOAT_CAST(pParams, Substrate.B_sp);
-      FLOAT_CAST(pParams, Substrate.Eth_ie);
-      FLOAT_CAST(pParams, Substrate.Eth_sp);
-      FLOAT_CAST(pParams, Substrate.k_sigma);
-      FLOAT_CAST(pParams, Substrate.beta_sigma);
-      FLOAT_CAST(pParams, Substrate.rho);
-
-      FLOAT_CAST(pParams, Passivation.A_ie);
-      FLOAT_CAST(pParams, Passivation.Eth_ie);
-
-      FLOAT_CAST(pParams, Ions.exponent);
-      FLOAT_CAST(pParams, Ions.meanEnergy);
-      FLOAT_CAST(pParams, Ions.sigmaEnergy);
-      FLOAT_CAST(pParams, Ions.inflectAngle);
-      FLOAT_CAST(pParams, Ions.n_l);
-      FLOAT_CAST(pParams, Ions.minAngle);
-      FLOAT_CAST(pParams, Ions.thetaRMin);
-      FLOAT_CAST(pParams, Ions.thetaRMax);
-    }
-    return pParams;
+#ifdef VIENNACORE_COMPILE_GPU
+struct PlasmaEtchingParametersGPU {
+  PlasmaEtchingParametersGPU() = default;
+  template <typename NumericType>
+  PlasmaEtchingParametersGPU(
+      const PlasmaEtchingParameters<NumericType> &parameters) {
+    Mask.A_sp = static_cast<float>(parameters.Mask.A_sp);
+    Mask.B_sp = static_cast<float>(parameters.Mask.B_sp);
+    Mask.Eth_sp = static_cast<float>(parameters.Mask.Eth_sp);
+    Mask.rho = static_cast<float>(parameters.Mask.rho);
+
+    Polymer.A_sp = static_cast<float>(parameters.Polymer.A_sp);
+    Polymer.B_sp = static_cast<float>(parameters.Polymer.B_sp);
+    Polymer.Eth_sp = static_cast<float>(parameters.Polymer.Eth_sp);
+    Polymer.rho = static_cast<float>(parameters.Polymer.rho);
+
+    Substrate.rho = static_cast<float>(parameters.Substrate.rho);
+    Substrate.Eth_sp = static_cast<float>(parameters.Substrate.Eth_sp);
+    Substrate.Eth_ie = static_cast<float>(parameters.Substrate.Eth_ie);
+    Substrate.A_sp = static_cast<float>(parameters.Substrate.A_sp);
+    Substrate.B_sp = static_cast<float>(parameters.Substrate.B_sp);
+    Substrate.A_ie = static_cast<float>(parameters.Substrate.A_ie);
+    Substrate.B_ie = static_cast<float>(parameters.Substrate.B_ie);
+    Substrate.k_sigma = static_cast<float>(parameters.Substrate.k_sigma);
+    Substrate.beta_sigma = static_cast<float>(parameters.Substrate.beta_sigma);
+
+    Passivation.Eth_ie = static_cast<float>(parameters.Passivation.Eth_ie);
+    Passivation.A_ie = static_cast<float>(parameters.Passivation.A_ie);
+
+    Ions.meanEnergy = static_cast<float>(parameters.Ions.meanEnergy);
+    Ions.sigmaEnergy = static_cast<float>(parameters.Ions.sigmaEnergy);
+    Ions.exponent = static_cast<float>(parameters.Ions.exponent);
+    Ions.inflectAngle = static_cast<float>(parameters.Ions.inflectAngle);
+    Ions.n_l = static_cast<float>(parameters.Ions.n_l);
+    Ions.minAngle = static_cast<float>(parameters.Ions.minAngle);
+    Ions.thetaRMin = static_cast<float>(parameters.Ions.thetaRMin);
+    Ions.thetaRMax = static_cast<float>(parameters.Ions.thetaRMax);
   }
+
+  PlasmaEtchingParameters<float>::MaskType Mask;
+  PlasmaEtchingParameters<float>::PolymerType Polymer;
+  PlasmaEtchingParameters<float>::MaterialType Substrate;
+  PlasmaEtchingParameters<float>::PassivationType Passivation;
+  PlasmaEtchingParameters<float>::IonType Ions;
 };
+#endif
 
 } // namespace viennaps
-
-#undef FLOAT_CAST