1 #include "sycldevice.hpp" 2 #include <CL/sycl.hpp> 3 4 namespace Petsc 5 { 6 7 namespace device 8 { 9 10 namespace sycl 11 { 12 13 namespace impl 14 { 15 16 class DeviceContext { 17 public: 18 struct PetscDeviceContext_IMPLS { 19 ::sycl::event event; 20 ::sycl::event begin; // timer-only 21 ::sycl::event end; // timer-only 22 #if PetscDefined(USE_DEBUG) 23 PetscBool timerInUse; 24 #endif 25 }; 26 27 private: 28 static bool initialized_; 29 30 static PetscErrorCode finalize_() noexcept 31 { 32 PetscFunctionBegin; 33 initialized_ = false; 34 PetscFunctionReturn(PETSC_SUCCESS); 35 } 36 37 static PetscErrorCode initialize_(PetscInt id, DeviceContext *dci) noexcept 38 { 39 PetscFunctionBegin; 40 PetscCall(PetscDeviceCheckDeviceCount_Internal(id)); 41 if (!initialized_) { 42 initialized_ = true; 43 PetscCall(PetscRegisterFinalize(finalize_)); 44 } 45 PetscFunctionReturn(PETSC_SUCCESS); 46 } 47 48 public: 49 const struct _DeviceContextOps ops = {destroy, changeStreamType, setUp, query, waitForContext, synchronize, getBlasHandle, getSolverHandle, getStreamHandle, beginTimer, endTimer, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr, nullptr}; 50 51 // default constructor 52 DeviceContext() noexcept = default; 53 54 // All of these functions MUST be static in order to be callable from C, otherwise they 55 // get the implicit 'this' pointer tacked on 56 static PetscErrorCode destroy(PetscDeviceContext dctx) noexcept 57 { 58 PetscFunctionBegin; 59 delete static_cast<PetscDeviceContext_IMPLS *>(dctx->data); 60 dctx->data = nullptr; 61 PetscFunctionReturn(PETSC_SUCCESS); 62 }; 63 64 static PetscErrorCode setUp(PetscDeviceContext dctx) noexcept 65 { 66 PetscFunctionBegin; 67 #if PetscDefined(USE_DEBUG) 68 static_cast<PetscDeviceContext_SYCL *>(dctx->data)->timerInUse = PETSC_FALSE; 69 #endif 70 // petsc/sycl currently only uses Kokkos's default execution space (and its queue), 71 // so in some sense, we have only one petsc device context. 72 PetscCall(PetscKokkosInitializeCheck()); 73 static_cast<PetscDeviceContext_SYCL *>(dctx->data)->queue = Kokkos::DefaultExecutionSpace().sycl_queue(); 74 PetscFunctionReturn(PETSC_SUCCESS); 75 }; 76 77 static PetscErrorCode query(PetscDeviceContext dctx, PetscBool *idle) noexcept 78 { 79 PetscFunctionBegin; 80 // available in future, https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/supported/sycl_ext_oneapi_queue_empty.asciidoc 81 // *idle = static_cast<PetscDeviceContext_SYCL*>(dctx->data)->queue.empty() ? PETSC_TRUE : PETSC_FALSE; 82 *idle = PETSC_FALSE; 83 PetscFunctionReturn(PETSC_SUCCESS); 84 }; 85 86 static PetscErrorCode synchronize(PetscDeviceContext dctx) noexcept 87 { 88 PetscBool idle = PETSC_TRUE; 89 const auto dci = static_cast<PetscDeviceContext_SYCL *>(dctx->data); 90 91 PetscFunctionBegin; 92 PetscCall(query(dctx, &idle)); 93 if (!idle) PetscCallCXX(dci->queue.wait()); 94 PetscFunctionReturn(PETSC_SUCCESS); 95 }; 96 97 static PetscErrorCode getStreamHandle(PetscDeviceContext dctx, void *handle) noexcept 98 { 99 PetscFunctionBegin; 100 *static_cast<::sycl::queue **>(handle) = &(static_cast<PetscDeviceContext_SYCL *>(dctx->data)->queue); 101 PetscFunctionReturn(PETSC_SUCCESS); 102 }; 103 104 static PetscErrorCode beginTimer(PetscDeviceContext dctx) noexcept 105 { 106 const auto dci = static_cast<PetscDeviceContext_SYCL *>(dctx->data); 107 108 PetscFunctionBegin; 109 #if PetscDefined(USE_DEBUG) 110 PetscCheck(!dci->timerInUse, PETSC_COMM_SELF, PETSC_ERR_PLIB, "Forgot to call PetscLogGpuTimeEnd()?"); 111 dci->timerInUse = PETSC_TRUE; 112 #endif 113 PetscCallCXX(dci->timeBegin = dci->timer.seconds()); 114 PetscFunctionReturn(PETSC_SUCCESS); 115 }; 116 117 static PetscErrorCode endTimer(PetscDeviceContext dctx, PetscLogDouble *elapsed) noexcept 118 { 119 const auto dci = static_cast<PetscDeviceContext_SYCL *>(dctx->data); 120 121 PetscFunctionBegin; 122 #if PetscDefined(USE_DEBUG) 123 PetscCheck(dci->timerInUse, PETSC_COMM_SELF, PETSC_ERR_PLIB, "Forgot to call PetscLogGpuTimeBegin()?"); 124 dci->timerInUse = PETSC_FALSE; 125 #endif 126 PetscCallCXX(dci->queue.wait()); 127 PetscCallCXX(*elapsed = dci->timer.seconds() - dci->timeBegin); 128 PetscFunctionReturn(PETSC_SUCCESS); 129 }; 130 131 static PetscErrorCode changeStreamType(PetscDeviceContext, PetscStreamType) noexcept { SETERRQ(PETSC_COMM_SELF, PETSC_ERR_SUP, "Not implemented"); }; 132 static PetscErrorCode setUp(PetscDeviceContext) noexcept { return PETSC_SUCCESS; }; // Nothing to setup 133 static PetscErrorCode query(PetscDeviceContext, PetscBool *) noexcept { SETERRQ(PETSC_COMM_SELF, PETSC_ERR_SUP, "Not implemented"); }; 134 static PetscErrorCode waitForContext(PetscDeviceContext, PetscDeviceContext) noexcept { SETERRQ(PETSC_COMM_SELF, PETSC_ERR_SUP, "Not implemented"); }; 135 static PetscErrorCode synchronize(PetscDeviceContext) noexcept { SETERRQ(PETSC_COMM_SELF, PETSC_ERR_SUP, "Not implemented"); }; 136 static PetscErrorCode getBlasHandle(PetscDeviceContext, void *) noexcept { SETERRQ(PETSC_COMM_SELF, PETSC_ERR_SUP, "Not implemented"); }; 137 static PetscErrorCode getSolverHandle(PetscDeviceContext, void *) noexcept { SETERRQ(PETSC_COMM_SELF, PETSC_ERR_SUP, "Not implemented"); }; 138 static PetscErrorCode getStreamHandle(PetscDeviceContext, void *) noexcept { SETERRQ(PETSC_COMM_SELF, PETSC_ERR_SUP, "Not implemented"); }; 139 static PetscErrorCode beginTimer(PetscDeviceContext) noexcept { SETERRQ(PETSC_COMM_SELF, PETSC_ERR_SUP, "Not implemented"); }; 140 static PetscErrorCode endTimer(PetscDeviceContext, PetscLogDouble *) noexcept { SETERRQ(PETSC_COMM_SELF, PETSC_ERR_SUP, "Not implemented"); }; 141 }; 142 143 } // namespace impl 144 145 } // namespace sycl 146 147 } // namespace device 148 149 } // namespace Petsc 150 151 PetscErrorCode PetscDeviceContextCreate_SYCL(PetscDeviceContext dctx) 152 { 153 using namespace Petsc::device::sycl::impl; 154 155 static const DeviceContext syclctx; 156 157 PetscFunctionBegin; 158 dctx->data = new DeviceContext::PetscDeviceContext_IMPLS(); 159 PetscCall(PetscMemcpy(dctx->ops, &syclctx.ops, sizeof(syclctx.ops))); 160 PetscFunctionReturn(PETSC_SUCCESS); 161 } 162