Method addTimeIntegratedPointSources

src/Kernels/PointSourceClusterOnDevice.cpp:32–102

unsigned PointSourceClusterOnDevice::size() const { return sources_->numberOfSources; }

void PointSourceClusterOnDevice::addTimeIntegratedPointSources(
    double from, double to, seissol::parallel::runtime::StreamRuntime& runtime) {
  auto& queue = seissol::AcceleratorDevice::getInstance().getSyclDefaultQueue();
  auto& mapping = clusterMapping_->cellToSources;
  if (mapping.size() > 0) {
    runtime.syncToSycl(&queue);

    auto* mappingPtr = mapping.data();
    auto* mInvJInvPhisAtSources = sources_->mInvJInvPhisAtSources.data();
    auto* tensor = sources_->tensor.data();
    auto* a = sources_->A.data();
    auto* stiffnessTensor = sources_->stiffnessTensor.data();
    auto* onsetTime = sources_->onsetTime.data();
    auto* samplingInterval = sources_->samplingInterval.data();
    auto sampleOffsets = std::array<std::size_t*, 3u>{sources_->sampleOffsets[0].data(),
                                                      sources_->sampleOffsets[1].data(),
                                                      sources_->sampleOffsets[2].data()};
    auto sample = std::array<real*, 3u>{
        sources_->sample[0].data(), sources_->sample[1].data(), sources_->sample[2].data()};

    sycl::range rng{mapping.size()};
    if (sources_->mode == sourceterm::PointSourceMode::Nrf) {
      queue.submit([&](sycl::handler& cgh) {
        cgh.parallel_for(rng, [=](sycl::item<1> id) {
          const unsigned startSource = mappingPtr[id[0]].pointSourcesOffset;
          const unsigned endSource =
              mappingPtr[id[0]].pointSourcesOffset + mappingPtr[id[0]].numberOfPointSources;
          for (unsigned source = startSource; source < endSource; ++source) {
            std::array<real, 3u> slip;
            for (int i = 0; i < 3; ++i) {
              auto o0 = sampleOffsets[i][source];
              auto o1 = sampleOffsets[i][source + 1];
              slip[i] = computeSampleTimeIntegral<seissol::functions::SyclStdFunctions>(
                  from, to, onsetTime[source], samplingInterval[source], sample[i] + o0, o1 - o0);
            }

            addTimeIntegratedPointSourceNRF(slip,
                                            mInvJInvPhisAtSources[source].data(),
                                            tensor[source].data(),
                                            a[source],
                                            stiffnessTensor[source].data(),
                                            from,
                                            to,
                                            *mappingPtr[id[0]].dofs);
          }
        });
      });
    } else {
      queue.submit([&](sycl::handler& cgh) {
        cgh.parallel_for(rng, [=](sycl::item<1> id) {
          const unsigned startSource = mappingPtr[id[0]].pointSourcesOffset;
          const unsigned endSource =
              mappingPtr[id[0]].pointSourcesOffset + mappingPtr[id[0]].numberOfPointSources;
          for (unsigned source = startSource; source < endSource; ++source) {
            auto o0 = sampleOffsets[0][source];
            auto o1 = sampleOffsets[0][source + 1];
            const real slip = computeSampleTimeIntegral<seissol::functions::SyclStdFunctions>(
                from, to, onsetTime[source], samplingInterval[source], sample[0] + o0, o1 - o0);