diff --git a/src/shared/particle_dynamics/dynamics_algorithms.h b/src/shared/particle_dynamics/dynamics_algorithms.h index d1506004b4..71f514b64b 100644 --- a/src/shared/particle_dynamics/dynamics_algorithms.h +++ b/src/shared/particle_dynamics/dynamics_algorithms.h @@ -141,10 +141,11 @@ class SimpleDynamicsCK : public LocalDynamicsType, public BaseDynamics { this->setUpdated(this->identifier_.getSPHBody()); this->setupDynamics(dt); - particle_for(kernel_implementation_, + ComputingKernel *computing_kernel = kernel_implementation_.getComputingKernel(); + particle_for(ExecutionPolicy{}, this->identifier_.LoopRange(), - [=](size_t i, auto &&computing_kernel) - { computing_kernel.update(i, dt); }); + [=](size_t i) + { computing_kernel->update(i, dt); }); }; protected: diff --git a/src/shared/particle_dynamics/execution/execution.h b/src/shared/particle_dynamics/execution/execution.h index 70baa55e9f..097cee6fb6 100644 --- a/src/shared/particle_dynamics/execution/execution.h +++ b/src/shared/particle_dynamics/execution/execution.h @@ -50,21 +50,21 @@ class Implementation public: explicit Implementation(LocalDynamicsType &local_dynamics) - : local_dynamics_(local_dynamics), delegated_kernel_(nullptr) {} + : local_dynamics_(local_dynamics), computing_kernel_(nullptr) {} - ComputingKernel *getDelegatedKernel() + ComputingKernel *getComputingKernel() { - if (delegated_kernel_ == nullptr) + if (computing_kernel_ == nullptr) { - delegated_kernel_ = + computing_kernel_ = kernel_ptr_keeper_.template createPtr(local_dynamics_); } - return delegated_kernel_; + return computing_kernel_; } private: LocalDynamicsType &local_dynamics_; - ComputingKernel *delegated_kernel_; + ComputingKernel *computing_kernel_; }; } // namespace execution } // namespace SPH diff --git a/src/shared/particle_dynamics/particle_iterators.h b/src/shared/particle_dynamics/particle_iterators.h index 1ce5099133..3a1a2760ee 100644 --- a/src/shared/particle_dynamics/particle_iterators.h +++ b/src/shared/particle_dynamics/particle_iterators.h @@ -74,32 +74,6 @@ inline void particle_for(const ParallelPolicy &par, const IndexRange &particles_ ap); }; -template -inline void particle_for(Implementation &kernel_implementation, - const IndexRange &particles_range, const ComputingKernelFunction &kernel_function) -{ - auto delegated_kernel = kernel_implementation.getDelegatedKernel(); - for (size_t i = particles_range.begin(); i < particles_range.end(); ++i) - kernel_function(i, *delegated_kernel); -}; - -template -inline void particle_for(Implementation &kernel_implementation, - const IndexRange &particles_range, const ComputingKernelFunction &kernel_function) -{ - auto delegated_kernel = kernel_implementation.getDelegatedKernel(); - parallel_for( - particles_range, - [&](const IndexRange &r) - { - for (size_t i = r.begin(); i < r.end(); ++i) - { - kernel_function(i, *delegated_kernel); - } - }, - ap); -}; - /** * Bodypart By Particle-wise iterators (for sequential and parallel computing). */ diff --git a/src/src_sycl/shared/particle_dynamics/execution_sycl.h b/src/src_sycl/shared/particle_dynamics/execution_sycl.h index 55355413f7..987b6b56b2 100644 --- a/src/src_sycl/shared/particle_dynamics/execution_sycl.h +++ b/src/src_sycl/shared/particle_dynamics/execution_sycl.h @@ -115,37 +115,69 @@ class ExecutionEvent private: std::vector event_list_; }; +} // namespace execution + +/* SYCL memory transfer utilities */ +template +inline T *allocateDeviceOnly(std::size_t size) +{ + return sycl::malloc_device(size, execution::execution_instance.getQueue()); +} + +template +inline T *allocateDeviceShared(std::size_t size) +{ + return sycl::malloc_shared(size, execution::execution_instance.getQueue()); +} + +template +inline void freeDeviceData(T *device_mem) +{ + sycl::free(device_mem, execution::execution_instance.getQueue()); +} + +template +inline execution::ExecutionEvent copyToDevice(const T *host, T *device, std::size_t size) +{ + return execution::execution_instance.getQueue().memcpy(device, host, size * sizeof(T)); +} +template +inline execution::ExecutionEvent copyFromDevice(T *host, const T *device, std::size_t size) +{ + return execution::execution_instance.getQueue().memcpy(host, device, size * sizeof(T)); +} + +namespace execution +{ template class Implementation { using ComputingKernel = typename LocalDynamicsType::ComputingKernel; - using ComputingKernelBuffer = sycl::buffer; - UniquePtrKeeper kernel_ptr_keeper_; - UniquePtrKeeper kernel_buffer_ptr_keeper_; public: explicit Implementation(LocalDynamicsType &local_dynamics) - : local_dynamics_(local_dynamics), computing_kernel_(nullptr), - computing_kernel_buffer_(nullptr) {} - - ComputingKernelBuffer &getBuffer() + : local_dynamics_(local_dynamics), computing_kernel_(nullptr) {} + ~Implementation() + { + freeDeviceData(computing_kernel_); + } + ComputingKernel *getComputingKernel() { if (computing_kernel_ == nullptr) { - computing_kernel_ = kernel_ptr_keeper_ - .template createPtr(local_dynamics_); - computing_kernel_buffer_ = kernel_buffer_ptr_keeper_ - .template createPtr(computing_kernel_, 1); + + computing_kernel_ = allocateDeviceOnly(1); + ComputingKernel host = ComputingKernel(local_dynamics_); + copyToDevice(&host, computing_kernel_, 1); } - return *computing_kernel_buffer_; + return computing_kernel_; } private: LocalDynamicsType &local_dynamics_; ComputingKernel *computing_kernel_; - ComputingKernelBuffer *computing_kernel_buffer_; }; } // namespace execution } // namespace SPH diff --git a/src/src_sycl/shared/particle_dynamics/particle_iterators_sycl.h b/src/src_sycl/shared/particle_dynamics/particle_iterators_sycl.h index 7080ec5178..adbcbe73b3 100644 --- a/src/src_sycl/shared/particle_dynamics/particle_iterators_sycl.h +++ b/src/src_sycl/shared/particle_dynamics/particle_iterators_sycl.h @@ -34,20 +34,17 @@ namespace SPH { -template -inline void particle_for(Implementation &kernel_implementation, - const IndexRange &particles_range, const ComputingKernelFunction &kernel_function) +template +inline void particle_for(const ParallelDevicePolicy &par_device, + const IndexRange &particles_range, const LocalDynamicsFunction &local_dynamics_function) { auto &sycl_queue = execution_instance.getQueue(); - auto &sycl_buffer = kernel_implementation.getBuffer(); const size_t particles_size = particles_range.size(); sycl_queue.submit([&](sycl::handler &cgh) - { - auto sycl_accessor = sycl_buffer.get_access(cgh, sycl::read_write); - cgh.parallel_for(execution_instance.getUniformNdRange(particles_size), [=](sycl::nd_item<1> index) { + { cgh.parallel_for(execution_instance.getUniformNdRange(particles_size), [=](sycl::nd_item<1> index) + { if(index.get_global_id(0) < particles_size) - kernel_function(index.get_global_id(0), sycl_accessor[0]); - }); }) + local_dynamics_function(index.get_global_id(0)); }); }) .wait_and_throw(); } } // namespace SPH diff --git a/src/src_sycl/shared/variable/base_variable_sycl.h b/src/src_sycl/shared/variable/base_variable_sycl.h deleted file mode 100644 index 7a3c4230e4..0000000000 --- a/src/src_sycl/shared/variable/base_variable_sycl.h +++ /dev/null @@ -1,67 +0,0 @@ -/* ------------------------------------------------------------------------- * - * SPHinXsys * - * ------------------------------------------------------------------------- * - * SPHinXsys (pronunciation: s'finksis) is an acronym from Smoothed Particle * - * Hydrodynamics for industrial compleX systems. It provides C++ APIs for * - * physical accurate simulation and aims to model coupled industrial dynamic * - * systems including fluid, solid, multi-body dynamics and beyond with SPH * - * (smoothed particle hydrodynamics), a meshless computational method using * - * particle discretization. * - * * - * SPHinXsys is partially funded by German Research Foundation * - * (Deutsche Forschungsgemeinschaft) DFG HU1527/6-1, HU1527/10-1, * - * HU1527/12-1 and HU1527/12-4. * - * * - * Portions copyright (c) 2017-2023 Technical University of Munich and * - * the authors' affiliations. * - * * - * Licensed under the Apache License, Version 2.0 (the "License"); you may * - * not use this file except in compliance with the License. You may obtain a * - * copy of the License at http://www.apache.org/licenses/LICENSE-2.0. * - * * - * ------------------------------------------------------------------------- */ -/** - * @file base_variable_sycl.h - * @brief TBD. - * @author Alberto Guarnieri and Xiangyu Hu - */ -#ifndef BASE_VARIABLE_SYCL_H -#define BASE_VARIABLE_SYCL_H - -#include "execution_sycl.h" - -namespace SPH -{ -/* SYCL memory transfer utilities */ -template -inline T *allocateDeviceOnly(std::size_t size) -{ - return sycl::malloc_device(size, execution::execution_instance.getQueue()); -} - -template -inline T *allocateDeviceShared(std::size_t size) -{ - return sycl::malloc_shared(size, execution::execution_instance.getQueue()); -} - -template -inline void freeDeviceData(T *device_mem) -{ - sycl::free(device_mem, execution::execution_instance.getQueue()); -} - -template -inline execution::ExecutionEvent copyToDevice(const T *host, T *device, std::size_t size) -{ - return execution::execution_instance.getQueue().memcpy(device, host, size * sizeof(T)); -} - -template -inline execution::ExecutionEvent copyFromDevice(T *host, const T *device, std::size_t size) -{ - return execution::execution_instance.getQueue().memcpy(host, device, size * sizeof(T)); -} -} // namespace SPH - -#endif // BASE_VARIABLE_SYCL_H \ No newline at end of file diff --git a/src/src_sycl/shared/variable/base_variable_sycl.hpp b/src/src_sycl/shared/variable/base_variable_sycl.hpp index a91c150e6a..09dd6f0a18 100644 --- a/src/src_sycl/shared/variable/base_variable_sycl.hpp +++ b/src/src_sycl/shared/variable/base_variable_sycl.hpp @@ -29,7 +29,7 @@ #define BASE_VARIABLE_SYCL_HPP #include "base_variable.h" -#include "base_variable_sycl.h" +#include "execution_sycl.h" namespace SPH {