/*
* This file is part of the GROMACS molecular simulation package.
*
- * Copyright (c) 2020, by the GROMACS development team, led by
+ * Copyright (c) 2020,2021, by the GROMACS development team, led by
* Mark Abraham, David van der Spoel, Berk Hess, and Erik Lindahl,
* and including many others, as listed in the AUTHORS file in the
* top-level source directory and at http://www.gromacs.org.
* \inlibraryapi
*/
+#include <utility>
+
#include "gromacs/gpu_utils/device_context.h"
#include "gromacs/gpu_utils/device_stream.h"
#include "gromacs/gpu_utils/devicebuffer_datatype.h"
#include "gromacs/gpu_utils/gmxsycl.h"
#include "gromacs/gpu_utils/gpu_utils.h" //only for GpuApiCallBehavior
#include "gromacs/gpu_utils/gputraits_sycl.h"
+#include "gromacs/utility/fatalerror.h"
#include "gromacs/utility/gmxassert.h"
#include "gromacs/utility/stringutil.h"
//! Copy constructor (references the same underlying SYCL buffer).
template<typename T>
-DeviceBuffer<T>::DeviceBuffer(DeviceBuffer<T> const& src) :
- buffer_(new ClSyclBufferWrapper(*src.buffer_))
+DeviceBuffer<T>::DeviceBuffer(DeviceBuffer<T> const& src)
{
+ if (src.buffer_)
+ {
+ buffer_ = std::make_unique<ClSyclBufferWrapper>(*src.buffer_);
+ }
+ else
+ {
+ buffer_ = nullptr;
+ }
}
//! Move constructor.
template<typename T>
DeviceBuffer<T>& DeviceBuffer<T>::operator=(DeviceBuffer<T> const& src)
{
- buffer_.reset(new ClSyclBufferWrapper(*src.buffer_));
+ if (src.buffer_)
+ {
+ buffer_ = std::make_unique<ClSyclBufferWrapper>(*src.buffer_);
+ }
+ else
+ {
+ buffer_.reset(nullptr);
+ }
return *this;
}
return *this;
}
+
+namespace gmx::internal
+{
+//! Shorthand alias to create a placeholder SYCL accessor with chosen data type and access mode.
+template<class T, cl::sycl::access::mode mode>
+using PlaceholderAccessor =
+ cl::sycl::accessor<T, 1, mode, cl::sycl::access::target::global_buffer, cl::sycl::access::placeholder::true_t>;
+} // namespace gmx::internal
+
+/** \brief
+ * Thin wrapper around placeholder accessor that allows implicit construction from \c DeviceBuffer.
+ *
+ * "Placeholder accessor" is an indicator of the intent to create an accessor for certain buffer
+ * of a certain type, that is not yet bound to a specific command group handler (device). Such
+ * accessors can be created outside SYCL kernels, which is helpful if we want to pass them as
+ * function arguments.
+ *
+ * \tparam T Type of buffer content.
+ * \tparam mode Access mode.
+ */
+template<class T, cl::sycl::access::mode mode>
+class DeviceAccessor : public gmx::internal::PlaceholderAccessor<T, mode>
+{
+public:
+ // Inherit all the constructors
+ using gmx::internal::PlaceholderAccessor<T, mode>::PlaceholderAccessor;
+ //! Construct Accessor from DeviceBuffer (must be initialized)
+ DeviceAccessor(DeviceBuffer<T>& buffer) :
+ gmx::internal::PlaceholderAccessor<T, mode>(getSyclBuffer(buffer))
+ {
+ }
+ //! Construct read-only Accessor from a const DeviceBuffer (must be initialized)
+ DeviceAccessor(const DeviceBuffer<T>& buffer) :
+ gmx::internal::PlaceholderAccessor<T, mode>(getSyclBuffer(const_cast<DeviceBuffer<T>&>(buffer)))
+ {
+ /* There were some discussions about making it possible to create read-only sycl::accessor
+ * from a const sycl::buffer (https://github.com/KhronosGroup/SYCL-Docs/issues/10), but
+ * it did not make it into the SYCL2020 standard. So, we have to use const_cast above */
+ /* Using static_assert to ensure that only mode::read accessors can be created from a
+ * const DeviceBuffer. static_assert provides better error messages than std::enable_if. */
+ static_assert(mode == cl::sycl::access::mode::read,
+ "Can not create non-read-only accessor from a const DeviceBuffer");
+ }
+ void bind(cl::sycl::handler& cgh) { cgh.require(*this); }
+
+private:
+ //! Helper function to get sycl:buffer object from DeviceBuffer wrapper, with a sanity check.
+ static inline cl::sycl::buffer<T, 1>& getSyclBuffer(DeviceBuffer<T>& buffer)
+ {
+ GMX_ASSERT(bool(buffer), "Trying to construct accessor from an uninitialized buffer");
+ return *buffer.buffer_;
+ }
+};
+
+namespace gmx::internal
+{
+//! A non-functional class that can be used instead of real accessors
+template<class T>
+struct NullAccessor
+{
+ NullAccessor(const DeviceBuffer<T>& /*buffer*/) {}
+ //! Allow casting to nullptr
+ constexpr operator std::nullptr_t() const { return nullptr; }
+ //! Placeholder implementation of \c cl::sycl::accessor::get_pointer.
+ T* get_pointer() const noexcept { return nullptr; }
+ void bind(cl::sycl::handler& /*cgh*/) { assert(false); }
+};
+} // namespace gmx::internal
+
+/** \brief
+ * Helper class to be used as function argument. Will either correspond to a device accessor, or an empty class.
+ *
+ * Example usage:
+ * \code
+ template <bool doFoo>
+ void getBarKernel(handler& cgh, OptionalAccessor<float, mode::read, doFoo> a_fooPrms)
+ {
+ if constexpr (doFoo)
+ cgh.require(a_fooPrms);
+ // Can only use a_fooPrms if doFoo == true
+ }
+
+ template <bool doFoo>
+ void callBar(DeviceBuffer<float> b_fooPrms)
+ {
+ // If doFoo is false, b_fooPrms will be ignored (can be not initialized).
+ // Otherwise, an accessor will be built (b_fooPrms must be a valid buffer).
+ auto kernel = getBarKernel<doFoo>(b_fooPrms);
+ // If the accessor in not enabled, anything can be passed as its ctor argument.
+ auto kernel2 = getBarKernel<false>(nullptr_t);
+ }
+ * \endcode
+ *
+ * \tparam T Data type of the underlying buffer
+ * \tparam mode Access mode of the accessor
+ * \tparam enabled Compile-time flag indicating whether we want to actually create an accessor.
+ */
+template<class T, cl::sycl::access::mode mode, bool enabled>
+using OptionalAccessor =
+ std::conditional_t<enabled, DeviceAccessor<T, mode>, gmx::internal::NullAccessor<T>>;
+
#endif // #ifndef DOXYGEN
+/*! \brief Check the validity of the device buffer.
+ *
+ * Checks if the buffer is valid and if its allocation is big enough.
+ *
+ * \param[in] buffer Device buffer to be checked.
+ * \param[in] requiredSize Number of elements that the buffer will have to accommodate.
+ *
+ * \returns Whether the device buffer exists and has enough capacity.
+ */
+template<typename T>
+static gmx_unused bool checkDeviceBuffer(const DeviceBuffer<T>& buffer, int requiredSize)
+{
+ return buffer.buffer_ && (static_cast<int>(buffer.buffer_->get_count()) >= requiredSize);
+}
+
/*! \libinternal \brief
* Allocates a device-side buffer.
* It is currently a caller's responsibility to call it only on not-yet allocated buffers.
return; // such calls are actually made with empty domains
}
GMX_ASSERT(buffer, "needs a buffer pointer");
- GMX_ASSERT(buffer->buffer_, "needs an initialized buffer pointer");
GMX_ASSERT(hostBuffer, "needs a host buffer pointer");
+ GMX_ASSERT(checkDeviceBuffer(*buffer, startingOffset + numValues),
+ "buffer too small or not initialized");
+
cl::sycl::buffer<ValueType>& syclBuffer = *buffer->buffer_;
cl::sycl::event ev = deviceStream.stream().submit([&](cl::sycl::handler& cgh) {
GMX_ASSERT(buffer, "needs a buffer pointer");
GMX_ASSERT(hostBuffer, "needs a host buffer pointer");
+ GMX_ASSERT(checkDeviceBuffer(*buffer, startingOffset + numValues),
+ "buffer too small or not initialized");
+
cl::sycl::buffer<ValueType>& syclBuffer = *buffer->buffer_;
cl::sycl::event ev = deviceStream.stream().submit([&](cl::sycl::handler& cgh) {
}
}
+/*! \brief
+ * Performs the device-to-device data copy, synchronous or asynchronously on request.
+ *
+ * \tparam ValueType Raw value type of the \p buffer.
+ */
+template<typename ValueType>
+void copyBetweenDeviceBuffers(DeviceBuffer<ValueType>* /* destinationDeviceBuffer */,
+ DeviceBuffer<ValueType>* /* sourceDeviceBuffer */,
+ size_t /* numValues */,
+ const DeviceStream& /* deviceStream */,
+ GpuApiCallBehavior /* transferKind */,
+ CommandEvent* /*timingEvent*/)
+{
+ // SYCL-TODO
+ gmx_fatal(FARGS, "D2D copy stub was called. Not yet implemented in SYCL.");
+}
+
+
+namespace gmx::internal
+{
+/*! \brief Helper function to clear device buffer.
+ *
+ * Not applicable to GROMACS's Float3 (a.k.a. gmx::RVec) and other custom types.
+ * From SYCL specs: "T must be a scalar value or a SYCL vector type."
+ */
+template<typename ValueType>
+cl::sycl::event fillSyclBufferWithNull(cl::sycl::buffer<ValueType, 1>& buffer,
+ size_t startingOffset,
+ size_t numValues,
+ cl::sycl::queue queue)
+{
+ using cl::sycl::access::mode;
+ const cl::sycl::range<1> range(numValues);
+ const cl::sycl::id<1> offset(startingOffset);
+ const ValueType pattern = ValueType(0); // SYCL vectors support initialization by scalar
+
+ return queue.submit([&](cl::sycl::handler& cgh) {
+ auto d_bufferAccessor =
+ cl::sycl::accessor<ValueType, 1, mode::discard_write>{ buffer, cgh, range, offset };
+ cgh.fill(d_bufferAccessor, pattern);
+ });
+}
+
+//! \brief Helper function to clear device buffer of type Float3.
+template<>
+inline cl::sycl::event fillSyclBufferWithNull(cl::sycl::buffer<Float3, 1>& buffer,
+ size_t startingOffset,
+ size_t numValues,
+ cl::sycl::queue queue)
+{
+ constexpr bool usingHipSycl =
+#ifdef __HIPSYCL__
+ true;
+#else
+ false;
+#endif
+
+ if constexpr (usingHipSycl)
+ {
+ // hipSYCL does not support reinterpret but allows using Float3 directly.
+ using cl::sycl::access::mode;
+ const cl::sycl::range<1> range(numValues);
+ const cl::sycl::id<1> offset(startingOffset);
+ const Float3 pattern{ 0, 0, 0 };
+
+ return queue.submit([&](cl::sycl::handler& cgh) {
+ auto d_bufferAccessor =
+ cl::sycl::accessor<Float3, 1, mode::discard_write>{ buffer, cgh, range, offset };
+ cgh.fill(d_bufferAccessor, pattern);
+ });
+ }
+ else // When not using hipSYCL, reinterpret as a flat float array
+ {
+#ifndef __HIPSYCL__
+ cl::sycl::buffer<float, 1> bufferAsFloat = buffer.reinterpret<float, 1>(buffer.get_count() * DIM);
+ return fillSyclBufferWithNull<float>(
+ bufferAsFloat, startingOffset * DIM, numValues * DIM, std::move(queue));
+#endif
+ }
+}
+
+} // namespace gmx::internal
+
/*! \brief
* Clears the device buffer asynchronously.
*
}
GMX_ASSERT(buffer, "needs a buffer pointer");
- const ValueType pattern{};
- cl::sycl::buffer<ValueType>& syclBuffer = *(buffer->buffer_);
+ GMX_ASSERT(checkDeviceBuffer(*buffer, startingOffset + numValues),
+ "buffer too small or not initialized");
- cl::sycl::event ev = deviceStream.stream().submit([&](cl::sycl::handler& cgh) {
- auto d_bufferAccessor = cl::sycl::accessor<ValueType, 1, cl::sycl::access::mode::discard_write>{
- syclBuffer, cgh, cl::sycl::range(numValues), cl::sycl::id(startingOffset)
- };
- cgh.fill(d_bufferAccessor, pattern);
- });
-}
+ cl::sycl::buffer<ValueType>& syclBuffer = *(buffer->buffer_);
-/*! \brief Check the validity of the device buffer.
- *
- * Checks if the buffer is valid and if its allocation is big enough.
- *
- * \param[in] buffer Device buffer to be checked.
- * \param[in] requiredSize Number of elements that the buffer will have to accommodate.
- *
- * \returns Whether the device buffer exists and has enough capacity.
- */
-template<typename T>
-static gmx_unused bool checkDeviceBuffer(DeviceBuffer<T> buffer, int requiredSize)
-{
- return buffer.buffer_ && (static_cast<int>(buffer.buffer_->get_count()) >= requiredSize);
+ gmx::internal::fillSyclBufferWithNull<ValueType>(
+ syclBuffer, startingOffset, numValues, deviceStream.stream());
}
/*! \brief Create a texture object for an array of type ValueType.
* \param[in,out] deviceBuffer Device buffer to store data in.
*/
template<typename ValueType>
-void destroyParamLookupTable(DeviceBuffer<ValueType>* deviceBuffer, DeviceTexture& /* deviceTexture */)
+void destroyParamLookupTable(DeviceBuffer<ValueType>* deviceBuffer, DeviceTexture* /* deviceTexture */)
{
- deviceBuffer->buffer_.reset(nullptr);
+ freeDeviceBuffer(deviceBuffer);
}
#endif // GMX_GPU_UTILS_DEVICEBUFFER_SYCL_H
biod.pnpi.spb.ru / alexxy / gromacs.git / blobdiff