1: #include "sycldevice.hpp"
  2: #include <CL/sycl.hpp>
  3: #include <Kokkos_Core.hpp>

  5: namespace Petsc
  6: {

  8: namespace device
  9: {

 11: namespace sycl
 12: {

 14: namespace impl
 15: {

 17: class DeviceContext {
 18: public:
 19:   struct PetscDeviceContext_SYCL {
 20:     ::sycl::event event;
 21:     ::sycl::event begin;   // timer-only
 22:     ::sycl::event end;     // timer-only
 23:     Kokkos::Timer timer{}; // use cpu time since sycl events are return value of queue submission and we have no infrastructure to store them
 24:     double        timeBegin{};
 25: #if PetscDefined(USE_DEBUG)
 26:     PetscBool timerInUse{};
 27: #endif
 28:     ::sycl::queue queue;
 29:   };

 31: private:
 32:   static bool initialized_;

 34:   static PetscErrorCode finalize_() noexcept
 35:   {
 36:     PetscFunctionBegin;
 37:     initialized_ = false;
 38:     PetscFunctionReturn(PETSC_SUCCESS);
 39:   }

 41:   static PetscErrorCode initialize_(PetscInt id, PetscDeviceContext dctx) noexcept
 42:   {
 43:     PetscFunctionBegin;
 44:     PetscCall(PetscDeviceCheckDeviceCount_Internal(id));
 45:     if (!initialized_) {
 46:       initialized_ = true;
 47:       PetscCall(PetscRegisterFinalize(finalize_));
 48:     }
 49:     PetscFunctionReturn(PETSC_SUCCESS);
 50:   }

 52: public:
 53:   const struct _DeviceContextOps ops = {destroy, changeStreamType, setUp, query, waitForContext, synchronize, getBlasHandle, getSolverHandle, getStreamHandle, beginTimer, endTimer, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr};

 55:   // default constructor
 56:   DeviceContext() noexcept = default;

 58:   // All of these functions MUST be static in order to be callable from C, otherwise they
 59:   // get the implicit 'this' pointer tacked on
 60:   static PetscErrorCode destroy(PetscDeviceContext dctx) noexcept
 61:   {
 62:     PetscFunctionBegin;
 63:     delete static_cast<PetscDeviceContext_SYCL *>(dctx->data);
 64:     dctx->data = nullptr;
 65:     PetscFunctionReturn(PETSC_SUCCESS);
 66:   };

 68:   static PetscErrorCode setUp(PetscDeviceContext dctx) noexcept
 69:   {
 70:     PetscFunctionBegin;
 71: #if PetscDefined(USE_DEBUG)
 72:     static_cast<PetscDeviceContext_SYCL *>(dctx->data)->timerInUse = PETSC_FALSE;
 73: #endif
 74:     // petsc/sycl currently only uses Kokkos's default execution space (and its queue),
 75:     // so in some sense, we have only one petsc device context.
 76:     PetscCall(PetscKokkosInitializeCheck());
 77:     static_cast<PetscDeviceContext_SYCL *>(dctx->data)->queue = Kokkos::DefaultExecutionSpace().sycl_queue();
 78:     PetscFunctionReturn(PETSC_SUCCESS);
 79:   };

 81:   static PetscErrorCode query(PetscDeviceContext dctx, PetscBool *idle) noexcept
 82:   {
 83:     PetscFunctionBegin;
 84:     // available in future, https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/supported/sycl_ext_oneapi_queue_empty.asciidoc
 85:     // *idle = static_cast<PetscDeviceContext_SYCL*>(dctx->data)->queue.empty() ? PETSC_TRUE : PETSC_FALSE;
 86:     *idle = PETSC_FALSE;
 87:     PetscFunctionReturn(PETSC_SUCCESS);
 88:   };

 90:   static PetscErrorCode synchronize(PetscDeviceContext dctx) noexcept
 91:   {
 92:     PetscBool  idle = PETSC_TRUE;
 93:     const auto dci  = static_cast<PetscDeviceContext_SYCL *>(dctx->data);

 95:     PetscFunctionBegin;
 96:     PetscCall(query(dctx, &idle));
 97:     if (!idle) PetscCallCXX(dci->queue.wait());
 98:     PetscFunctionReturn(PETSC_SUCCESS);
 99:   };

101:   static PetscErrorCode getStreamHandle(PetscDeviceContext dctx, void **handle) noexcept
102:   {
103:     PetscFunctionBegin;
104:     *reinterpret_cast<::sycl::queue **>(handle) = &(static_cast<PetscDeviceContext_SYCL *>(dctx->data)->queue);
105:     PetscFunctionReturn(PETSC_SUCCESS);
106:   };

108:   static PetscErrorCode beginTimer(PetscDeviceContext dctx) noexcept
109:   {
110:     const auto dci = static_cast<PetscDeviceContext_SYCL *>(dctx->data);

112:     PetscFunctionBegin;
113: #if PetscDefined(USE_DEBUG)
114:     PetscCheck(!dci->timerInUse, PETSC_COMM_SELF, PETSC_ERR_PLIB, "Forgot to call PetscLogGpuTimeEnd()?");
115:     dci->timerInUse = PETSC_TRUE;
116: #endif
117:     PetscCallCXX(dci->timeBegin = dci->timer.seconds());
118:     PetscFunctionReturn(PETSC_SUCCESS);
119:   };

121:   static PetscErrorCode endTimer(PetscDeviceContext dctx, PetscLogDouble *elapsed) noexcept
122:   {
123:     const auto dci = static_cast<PetscDeviceContext_SYCL *>(dctx->data);

125:     PetscFunctionBegin;
126: #if PetscDefined(USE_DEBUG)
127:     PetscCheck(dci->timerInUse, PETSC_COMM_SELF, PETSC_ERR_PLIB, "Forgot to call PetscLogGpuTimeBegin()?");
128:     dci->timerInUse = PETSC_FALSE;
129: #endif
130:     PetscCallCXX(dci->queue.wait());
131:     PetscCallCXX(*elapsed = dci->timer.seconds() - dci->timeBegin);
132:     PetscFunctionReturn(PETSC_SUCCESS);
133:   };

135:   static PetscErrorCode changeStreamType(PetscDeviceContext, PetscStreamType) noexcept { SETERRQ(PETSC_COMM_SELF, PETSC_ERR_SUP, "Not implemented"); };
136:   static PetscErrorCode waitForContext(PetscDeviceContext, PetscDeviceContext) noexcept { SETERRQ(PETSC_COMM_SELF, PETSC_ERR_SUP, "Not implemented"); };
137:   static PetscErrorCode getBlasHandle(PetscDeviceContext, void *) noexcept { SETERRQ(PETSC_COMM_SELF, PETSC_ERR_SUP, "Not implemented"); };
138:   static PetscErrorCode getSolverHandle(PetscDeviceContext, void *) noexcept { SETERRQ(PETSC_COMM_SELF, PETSC_ERR_SUP, "Not implemented"); };
139: };

141: } // namespace impl

143: } // namespace sycl

145: } // namespace device

147: } // namespace Petsc

149: PetscErrorCode PetscDeviceContextCreate_SYCL(PetscDeviceContext dctx)
150: {
151:   using namespace Petsc::device::sycl::impl;

153:   static const DeviceContext syclctx;

155:   PetscFunctionBegin;
156:   PetscCallCXX(dctx->data = new DeviceContext::PetscDeviceContext_SYCL());
157:   dctx->ops[0] = syclctx.ops;
158:   PetscFunctionReturn(PETSC_SUCCESS);
159: }