1 /* 2 //@HEADER 3 // ************************************************************************ 4 // 5 // Kokkos v. 3.0 6 // Copyright (2020) National Technology & Engineering 7 // Solutions of Sandia, LLC (NTESS). 8 // 9 // Under the terms of Contract DE-NA0003525 with NTESS, 10 // the U.S. Government retains certain rights in this software. 11 // 12 // Redistribution and use in source and binary forms, with or without 13 // modification, are permitted provided that the following conditions are 14 // met: 15 // 16 // 1. Redistributions of source code must retain the above copyright 17 // notice, this list of conditions and the following disclaimer. 18 // 19 // 2. Redistributions in binary form must reproduce the above copyright 20 // notice, this list of conditions and the following disclaimer in the 21 // documentation and/or other materials provided with the distribution. 22 // 23 // 3. Neither the name of the Corporation nor the names of the 24 // contributors may be used to endorse or promote products derived from 25 // this software without specific prior written permission. 26 // 27 // THIS SOFTWARE IS PROVIDED BY NTESS "AS IS" AND ANY 28 // EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE 29 // IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR 30 // PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL NTESS OR THE 31 // CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, 32 // EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, 33 // PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR 34 // PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF 35 // LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING 36 // NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS 37 // SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. 38 // 39 // Questions? Contact Christian R. Trott (crtrott@sandia.gov) 40 // 41 // ************************************************************************ 42 //@HEADER 43 */ 44 45 #ifndef KOKKOS_SYCLSPACE_HPP 46 #define KOKKOS_SYCLSPACE_HPP 47 48 #include <Kokkos_Core_fwd.hpp> 49 50 #ifdef KOKKOS_ENABLE_SYCL 51 #include <Kokkos_Concepts.hpp> 52 #include <Kokkos_ScratchSpace.hpp> 53 #include <SYCL/Kokkos_SYCL_Instance.hpp> 54 #include <impl/Kokkos_SharedAlloc.hpp> 55 #include <impl/Kokkos_Tools.hpp> 56 57 namespace Kokkos { 58 namespace Experimental { 59 60 class SYCLDeviceUSMSpace { 61 public: 62 using execution_space = SYCL; 63 using memory_space = SYCLDeviceUSMSpace; 64 using device_type = Kokkos::Device<execution_space, memory_space>; 65 using size_type = Impl::SYCLInternal::size_type; 66 67 SYCLDeviceUSMSpace(); 68 explicit SYCLDeviceUSMSpace(sycl::queue queue); 69 70 void* allocate(const std::size_t arg_alloc_size) const; 71 void* allocate(const char* arg_label, const size_t arg_alloc_size, 72 const size_t arg_logical_size = 0) const; 73 74 void deallocate(void* const arg_alloc_ptr, 75 const std::size_t arg_alloc_size) const; 76 void deallocate(const char* arg_label, void* const arg_alloc_ptr, 77 const size_t arg_alloc_size, 78 const size_t arg_logical_size = 0) const; 79 80 private: 81 template <class, class, class, class> 82 friend class LogicalMemorySpace; 83 84 public: name()85 static constexpr const char* name() { return "SYCLDeviceUSM"; }; 86 87 private: 88 sycl::queue m_queue; 89 }; 90 91 class SYCLSharedUSMSpace { 92 public: 93 using execution_space = SYCL; 94 using memory_space = SYCLSharedUSMSpace; 95 using device_type = Kokkos::Device<execution_space, memory_space>; 96 using size_type = Impl::SYCLInternal::size_type; 97 98 SYCLSharedUSMSpace(); 99 explicit SYCLSharedUSMSpace(sycl::queue queue); 100 101 void* allocate(const std::size_t arg_alloc_size) const; 102 void* allocate(const char* arg_label, const size_t arg_alloc_size, 103 const size_t arg_logical_size = 0) const; 104 105 void deallocate(void* const arg_alloc_ptr, 106 const std::size_t arg_alloc_size) const; 107 void deallocate(const char* arg_label, void* const arg_alloc_ptr, 108 const size_t arg_alloc_size, 109 const size_t arg_logical_size = 0) const; 110 111 private: 112 template <class, class, class, class> 113 friend class LogicalMemorySpace; 114 115 public: name()116 static constexpr const char* name() { return "SYCLSharedUSM"; }; 117 118 private: 119 sycl::queue m_queue; 120 }; 121 } // namespace Experimental 122 123 namespace Impl { 124 static_assert(Kokkos::Impl::MemorySpaceAccess< 125 Kokkos::Experimental::SYCLDeviceUSMSpace, 126 Kokkos::Experimental::SYCLDeviceUSMSpace>::assignable, 127 ""); 128 129 static_assert(Kokkos::Impl::MemorySpaceAccess< 130 Kokkos::Experimental::SYCLSharedUSMSpace, 131 Kokkos::Experimental::SYCLSharedUSMSpace>::assignable, 132 ""); 133 134 template <> 135 struct MemorySpaceAccess<Kokkos::HostSpace, 136 Kokkos::Experimental::SYCLDeviceUSMSpace> { 137 enum : bool { assignable = false }; 138 enum : bool { accessible = false }; 139 enum : bool { deepcopy = true }; 140 }; 141 142 template <> 143 struct MemorySpaceAccess<Kokkos::HostSpace, 144 Kokkos::Experimental::SYCLSharedUSMSpace> { 145 // HostSpace::execution_space != SYCLSharedUSMSpace::execution_space 146 enum : bool { assignable = false }; 147 enum : bool { accessible = true }; 148 enum : bool { deepcopy = true }; 149 }; 150 151 template <> 152 struct MemorySpaceAccess<Kokkos::Experimental::SYCLDeviceUSMSpace, 153 Kokkos::HostSpace> { 154 enum : bool { assignable = false }; 155 enum : bool { accessible = false }; 156 enum : bool { deepcopy = true }; 157 }; 158 159 template <> 160 struct MemorySpaceAccess<Kokkos::Experimental::SYCLDeviceUSMSpace, 161 Kokkos::Experimental::SYCLSharedUSMSpace> { 162 // SYCLDeviceUSMSpace::execution_space == SYCLSharedUSMSpace::execution_space 163 enum : bool { assignable = true }; 164 enum : bool { accessible = true }; 165 enum : bool { deepcopy = true }; 166 }; 167 168 //---------------------------------------- 169 // SYCLSharedUSMSpace::execution_space == SYCL 170 // SYCLSharedUSMSpace accessible to both SYCL and Host 171 172 template <> 173 struct MemorySpaceAccess<Kokkos::Experimental::SYCLSharedUSMSpace, 174 Kokkos::HostSpace> { 175 enum : bool { assignable = false }; 176 enum : bool { accessible = false }; // SYCL cannot access HostSpace 177 enum : bool { deepcopy = true }; 178 }; 179 180 template <> 181 struct MemorySpaceAccess<Kokkos::Experimental::SYCLSharedUSMSpace, 182 Kokkos::Experimental::SYCLDeviceUSMSpace> { 183 // SYCLSharedUSMSpace::execution_space == SYCLDeviceUSMSpace::execution_space 184 // Can access SYCLSharedUSMSpace from Host but cannot access 185 // SYCLDeviceUSMSpace from Host 186 enum : bool { assignable = false }; 187 188 // SYCLSharedUSMSpace::execution_space can access SYCLDeviceUSMSpace 189 enum : bool { accessible = true }; 190 enum : bool { deepcopy = true }; 191 }; 192 193 template <> 194 struct MemorySpaceAccess< 195 Kokkos::Experimental::SYCLDeviceUSMSpace, 196 Kokkos::ScratchMemorySpace<Kokkos::Experimental::SYCL>> { 197 enum : bool { assignable = false }; 198 enum : bool { accessible = true }; 199 enum : bool { deepcopy = false }; 200 }; 201 202 template <> 203 struct MemorySpaceAccess< 204 Kokkos::Experimental::SYCLSharedUSMSpace, 205 Kokkos::ScratchMemorySpace<Kokkos::Experimental::SYCL>> { 206 enum : bool { assignable = false }; 207 enum : bool { accessible = true }; 208 enum : bool { deepcopy = false }; 209 }; 210 211 } // namespace Impl 212 213 namespace Impl { 214 215 template <> 216 class SharedAllocationRecord<Kokkos::Experimental::SYCLDeviceUSMSpace, void> 217 : public HostInaccessibleSharedAllocationRecordCommon< 218 Kokkos::Experimental::SYCLDeviceUSMSpace> { 219 private: 220 friend class SharedAllocationRecordCommon< 221 Kokkos::Experimental::SYCLDeviceUSMSpace>; 222 friend class HostInaccessibleSharedAllocationRecordCommon< 223 Kokkos::Experimental::SYCLDeviceUSMSpace>; 224 using base_t = HostInaccessibleSharedAllocationRecordCommon< 225 Kokkos::Experimental::SYCLDeviceUSMSpace>; 226 using RecordBase = SharedAllocationRecord<void, void>; 227 228 SharedAllocationRecord(const SharedAllocationRecord&) = delete; 229 SharedAllocationRecord(SharedAllocationRecord&&) = delete; 230 SharedAllocationRecord& operator=(const SharedAllocationRecord&) = delete; 231 SharedAllocationRecord& operator=(SharedAllocationRecord&&) = delete; 232 233 #ifdef KOKKOS_ENABLE_DEBUG 234 static RecordBase s_root_record; 235 #endif 236 237 const Kokkos::Experimental::SYCLDeviceUSMSpace m_space; 238 239 protected: 240 ~SharedAllocationRecord(); 241 242 SharedAllocationRecord( 243 const Kokkos::Experimental::SYCLDeviceUSMSpace& arg_space, 244 const std::string& arg_label, const size_t arg_alloc_size, 245 const RecordBase::function_type arg_dealloc = &base_t::deallocate); 246 }; 247 248 template <> 249 class SharedAllocationRecord<Kokkos::Experimental::SYCLSharedUSMSpace, void> 250 : public SharedAllocationRecordCommon< 251 Kokkos::Experimental::SYCLSharedUSMSpace> { 252 private: 253 friend class SharedAllocationRecordCommon< 254 Kokkos::Experimental::SYCLSharedUSMSpace>; 255 using base_t = 256 SharedAllocationRecordCommon<Kokkos::Experimental::SYCLSharedUSMSpace>; 257 using RecordBase = SharedAllocationRecord<void, void>; 258 259 SharedAllocationRecord(const SharedAllocationRecord&) = delete; 260 SharedAllocationRecord(SharedAllocationRecord&&) = delete; 261 SharedAllocationRecord& operator=(const SharedAllocationRecord&) = delete; 262 SharedAllocationRecord& operator=(SharedAllocationRecord&&) = delete; 263 264 static RecordBase s_root_record; 265 266 const Kokkos::Experimental::SYCLSharedUSMSpace m_space; 267 268 protected: 269 ~SharedAllocationRecord(); 270 271 SharedAllocationRecord() = default; 272 273 SharedAllocationRecord( 274 const Kokkos::Experimental::SYCLSharedUSMSpace& arg_space, 275 const std::string& arg_label, const size_t arg_alloc_size, 276 const RecordBase::function_type arg_dealloc = &base_t::deallocate); 277 }; 278 279 } // namespace Impl 280 281 } // namespace Kokkos 282 283 #endif 284 #endif 285