GCC Code Coverage Report

chrono_batch_render/
Coverage:
low: ≥ 0%
medium: ≥ 75.0%
high: ≥ 90.0%
Lines:
6 / 407
1.5%
Branches:
4 / 192
2.1%
Functions:
1 / 52
1.9%
device/vulkan/interop.cpp
Line Branch Exec Source
1 // SPDX-License-Identifier: BSD-3-Clause
2 #include "chrono_batch_render/device/vulkan/interop.h"
3
4 #include <condition_variable>
5 #include <cstring>
6 #include <deque>
7 #include <memory>
8 #include <mutex>
9 #include <thread>
10
11 #if !defined(_WIN32)
12 #include <unistd.h>
13 #endif
14
15 #include "chrono_batch_render/device/cuda/api.h"
16 #include "chrono_batch_render/device/cuda/device.h"
17 #include "chrono_batch_render/device/hip/api.h"
18 #include "chrono_batch_render/device/hip/device.h"
19 #include "chrono_batch_render/util/error.h"
20 #include "chrono_batch_render/util/log.h"
21
22 namespace chrono::batch_render::device::vulkan {
23
24 namespace {
25
26 constexpr uint32_t kVendorNvidia = 0x10de;
27 constexpr uint32_t kVendorAmd = 0x1002;
28
29 #if !defined(_WIN32)
30 ✗ int ExportFd(VulkanDevice& device, VkSemaphore semaphore) {
31 ✗ VkSemaphoreGetFdInfoKHR info{VK_STRUCTURE_TYPE_SEMAPHORE_GET_FD_INFO_KHR};
32 ✗ info.semaphore = semaphore;
33 ✗ info.handleType = VK_EXTERNAL_SEMAPHORE_HANDLE_TYPE_OPAQUE_FD_BIT;
34 ✗ int fd = -1;
35 ✗ Check(device.Functions().vkGetSemaphoreFdKHR(device.Handle(), &info, &fd),
36 "vkGetSemaphoreFdKHR");
37 ✗ return fd;
38 }
39 #endif
40
41 ✗ void CloseFd(int fd) {
42 #if !defined(_WIN32)
43 ✗ close(fd);
44 #else
45 (void)fd;
46 #endif
47 ✗ }
48
49 struct CudaDevice {
50 cuda::CUdevice device = 0;
51 int32_t ordinal = 0;
52 };
53
54 ✗ CudaDevice FindCudaDevice(const cuda::Api& cuda, const VulkanDevice& device) {
55 ✗ int count = 0;
56 ✗ cuda.Check(cuda.cuDeviceGetCount(&count), "cuDeviceGetCount");
57 ✗ CudaDevice found;
58
0/2
✗ Branch 0 not taken.
✗ Branch 1 not taken.
 ✗ for (int ordinal = 0; ordinal < count; ++ordinal) {
59 ✗ cuda::CUuuid uuid{};
60 ✗ cuda.Check(cuda.cuDeviceGet(&found.device, ordinal), "cuDeviceGet");
61 ✗ cuda.Check(cuda.cuDeviceGetUuid(&uuid, found.device), "cuDeviceGetUuid");
62
0/2
✗ Branch 0 not taken.
✗ Branch 1 not taken.
 ✗ if (std::memcmp(uuid.bytes, device.Adapter().uuid.data(), 16) == 0) {
63 ✗ found.ordinal = ordinal;
64 ✗ return found;
65 }
66 }
67
0/1
✗ Branch 2 not taken.
 ✗ throw UnsupportedError("CUDA has no GPU of the UUID of {}", device.Adapter().alias);
68 }
69
70 ✗ int FindHipDevice(const hip::Api& hip, const VulkanDevice& device) {
71 ✗ int count = 0;
72 ✗ hip.Check(hip.hipGetDeviceCount(&count), "hipGetDeviceCount");
73
0/2
✗ Branch 0 not taken.
✗ Branch 1 not taken.
 ✗ for (int ordinal = 0; ordinal < count; ++ordinal) {
74 ✗ hip::hipUUID uuid{};
75 ✗ hip.Check(hip.hipDeviceGetUuid(&uuid, ordinal), "hipDeviceGetUuid");
76
0/2
✗ Branch 0 not taken.
✗ Branch 1 not taken.
 ✗ if (std::memcmp(uuid.bytes, device.Adapter().uuid.data(), 16) == 0) {
77 ✗ return ordinal;
78 }
79 }
80
0/1
✗ Branch 2 not taken.
 ✗ throw UnsupportedError("HIP has no GPU of the UUID of {}", device.Adapter().alias);
81 }
82
83 class CudaContext {
84 public:
85 ✗ CudaContext(const cuda::Api& cuda, cuda::CUcontext context) : m_cuda(cuda) {
86 ✗ m_cuda.Check(m_cuda.cuCtxPushCurrent(context), "cuCtxPushCurrent");
87 ✗ }
88
89 ✗ ~CudaContext() {
90 ✗ cuda::CUcontext popped = nullptr;
91 ✗ m_cuda.Warn(m_cuda.cuCtxPopCurrent(&popped), "cuCtxPopCurrent");
92 ✗ }
93
94 CudaContext(const CudaContext&) = delete;
95 CudaContext& operator=(const CudaContext&) = delete;
96
97 private:
98 const cuda::Api& m_cuda;
99 };
100
101 #if !defined(_WIN32)
102 class CudaInterop : public StreamInterop {
103 public:
104 ✗ CudaInterop(VulkanDevice& device, VkSemaphore timeline, VkSemaphore follow)
105
0/1
✗ Branch 1 not taken.
 ✗ : m_cuda(cuda::GetApi()) {
106
0/1
✗ Branch 1 not taken.
 ✗ m_device = FindCudaDevice(m_cuda, device).device;
107
0/2
✗ Branch 1 not taken.
✗ Branch 4 not taken.
 ✗ m_cuda.Check(m_cuda.cuDevicePrimaryCtxRetain(&m_context, m_device),
108 "cuDevicePrimaryCtxRetain");
109
0/1
✗ Branch 1 not taken.
 ✗ const CudaContext scope(m_cuda, m_context);
110
0/1
✗ Branch 1 not taken.
 ✗ m_stream = cuda::Traits::CreateStream(m_cuda);
111
0/1
✗ Branch 1 not taken.
 ✗ m_timeline = Import(device, timeline);
112
0/1
✗ Branch 1 not taken.
 ✗ m_follow = Import(device, follow);
113 ✗ }
114
115 ✗ ~CudaInterop() override {
116 ✗ {
117 ✗ const CudaContext scope(m_cuda, m_context);
118
0/2
✗ Branch 0 not taken.
✗ Branch 1 not taken.
 ✗ if (m_follow) {
119 ✗ m_cuda.Warn(m_cuda.cuDestroyExternalSemaphore(m_follow),
120 "cuDestroyExternalSemaphore");
121 }
122
0/2
✗ Branch 0 not taken.
✗ Branch 1 not taken.
 ✗ if (m_timeline) {
123 ✗ m_cuda.Warn(m_cuda.cuDestroyExternalSemaphore(m_timeline),
124 "cuDestroyExternalSemaphore");
125 }
126
0/2
✗ Branch 0 not taken.
✗ Branch 1 not taken.
 ✗ if (m_stream) {
127 ✗ cuda::Traits::DestroyStream(m_cuda, m_stream);
128 }
129 ✗ }
130 ✗ m_cuda.Warn(m_cuda.cuDevicePrimaryCtxRelease(m_device), "cuDevicePrimaryCtxRelease");
131 ✗ }
132
133 ✗ void Follow(void* stream, uint64_t value) override {
134 ✗ const CudaContext scope(m_cuda, m_context);
135
0/1
✗ Branch 1 not taken.
 ✗ const cuda::CUevent event = cuda::Traits::CreateEvent(m_cuda, false);
136
0/2
✗ Branch 1 not taken.
✗ Branch 4 not taken.
 ✗ cuda::Traits::Record(m_cuda, event, cuda::Traits::ApplicationStream(stream), false);
137
0/1
✗ Branch 1 not taken.
 ✗ cuda::Traits::WaitEvent(m_cuda, m_stream, event);
138
0/1
✗ Branch 1 not taken.
 ✗ cuda::Traits::DestroyEvent(m_cuda, event);
139 ✗ cuda::CUDA_EXTERNAL_SEMAPHORE_SIGNAL_PARAMS params{};
140 ✗ params.params.fence.value = value;
141
0/2
✗ Branch 1 not taken.
✗ Branch 4 not taken.
 ✗ m_cuda.Check(m_cuda.cuSignalExternalSemaphoresAsync(&m_follow, &params, 1, m_stream),
142 "cuSignalExternalSemaphoresAsync");
143 ✗ }
144
145 ✗ void Wait(void* stream, uint64_t value) override {
146 ✗ const CudaContext scope(m_cuda, m_context);
147 ✗ cuda::CUDA_EXTERNAL_SEMAPHORE_WAIT_PARAMS params{};
148 ✗ params.params.fence.value = value;
149
0/3
✗ Branch 1 not taken.
✗ Branch 4 not taken.
✗ Branch 7 not taken.
 ✗ m_cuda.Check(m_cuda.cuWaitExternalSemaphoresAsync(&m_timeline, &params, 1,
150 cuda::Traits::ApplicationStream(stream)),
151 "cuWaitExternalSemaphoresAsync");
152 ✗ }
153
154 private:
155 ✗ cuda::CUexternalSemaphore Import(VulkanDevice& device, VkSemaphore semaphore) {
156 ✗ const int fd = ExportFd(device, semaphore);
157 ✗ cuda::CUDA_EXTERNAL_SEMAPHORE_HANDLE_DESC desc{};
158 ✗ desc.type = cuda::CU_EXTERNAL_SEMAPHORE_HANDLE_TYPE_TIMELINE_SEMAPHORE_FD;
159 ✗ desc.handle.fd = fd;
160 ✗ cuda::CUexternalSemaphore imported = nullptr;
161 ✗ const cuda::CUresult result = m_cuda.cuImportExternalSemaphore(&imported, &desc);
162
0/2
✗ Branch 0 not taken.
✗ Branch 1 not taken.
 ✗ if (result != 0) {
163 ✗ CloseFd(fd);
164 }
165 ✗ m_cuda.Check(result, "cuImportExternalSemaphore");
166 ✗ return imported;
167 }
168
169 const cuda::Api& m_cuda;
170 cuda::CUdevice m_device = 0;
171 cuda::CUcontext m_context = nullptr;
172 cuda::CUstream m_stream = nullptr;
173 cuda::CUexternalSemaphore m_timeline = nullptr;
174 cuda::CUexternalSemaphore m_follow = nullptr;
175 };
176
177 class HipInterop : public StreamInterop {
178 public:
179 ✗ HipInterop(VulkanDevice& device, VkSemaphore timeline, VkSemaphore follow)
180
0/1
✗ Branch 1 not taken.
 ✗ : m_hip(hip::GetApi()) {
181
0/1
✗ Branch 1 not taken.
 ✗ m_device = FindHipDevice(m_hip, device);
182
0/1
✗ Branch 1 not taken.
 ✗ Enter();
183
0/1
✗ Branch 1 not taken.
 ✗ m_stream = hip::Traits::CreateStream(m_hip);
184
0/1
✗ Branch 1 not taken.
 ✗ m_timeline = Import(device, timeline);
185
0/1
✗ Branch 1 not taken.
 ✗ m_follow = Import(device, follow);
186 ✗ }
187
188 ✗ ~HipInterop() override {
189 ✗ m_hip.Warn(m_hip.hipSetDevice(m_device), "hipSetDevice");
190
0/2
✗ Branch 0 not taken.
✗ Branch 1 not taken.
 ✗ if (m_follow) {
191 ✗ m_hip.Warn(m_hip.hipDestroyExternalSemaphore(m_follow), "hipDestroyExternalSemaphore");
192 }
193
0/2
✗ Branch 0 not taken.
✗ Branch 1 not taken.
 ✗ if (m_timeline) {
194 ✗ m_hip.Warn(m_hip.hipDestroyExternalSemaphore(m_timeline),
195 "hipDestroyExternalSemaphore");
196 }
197
0/2
✗ Branch 0 not taken.
✗ Branch 1 not taken.
 ✗ if (m_stream) {
198 ✗ hip::Traits::DestroyStream(m_hip, m_stream);
199 }
200 ✗ }
201
202 ✗ void Follow(void* stream, uint64_t value) override {
203 ✗ Enter();
204 ✗ const hip::hipEvent_t event = hip::Traits::CreateEvent(m_hip, false);
205 ✗ hip::Traits::Record(m_hip, event, hip::Traits::ApplicationStream(stream), false);
206 ✗ hip::Traits::WaitEvent(m_hip, m_stream, event);
207 ✗ hip::Traits::DestroyEvent(m_hip, event);
208 ✗ hip::hipExternalSemaphoreSignalParams params{};
209 ✗ params.params.fence.value = value;
210 ✗ m_hip.Check(m_hip.hipSignalExternalSemaphoresAsync(&m_follow, &params, 1, m_stream),
211 "hipSignalExternalSemaphoresAsync");
212 ✗ }
213
214 ✗ void Wait(void* stream, uint64_t value) override {
215 ✗ Enter();
216 ✗ hip::hipExternalSemaphoreWaitParams params{};
217 ✗ params.params.fence.value = value;
218 ✗ m_hip.Check(m_hip.hipWaitExternalSemaphoresAsync(&m_timeline, &params, 1,
219 hip::Traits::ApplicationStream(stream)),
220 "hipWaitExternalSemaphoresAsync");
221 ✗ }
222
223 private:
224 ✗ void Enter() const {
225 ✗ m_hip.Check(m_hip.hipSetDevice(m_device), "hipSetDevice");
226 ✗ }
227
228 ✗ hip::hipExternalSemaphore_t Import(VulkanDevice& device, VkSemaphore semaphore) {
229 ✗ const int fd = ExportFd(device, semaphore);
230 ✗ hip::hipExternalSemaphoreHandleDesc desc{};
231 ✗ desc.type = hip::hipExternalSemaphoreHandleTypeTimelineSemaphoreFd;
232 ✗ desc.handle.fd = fd;
233 ✗ hip::hipExternalSemaphore_t imported = nullptr;
234 ✗ const hip::hipError_t result = m_hip.hipImportExternalSemaphore(&imported, &desc);
235
0/2
✗ Branch 0 not taken.
✗ Branch 1 not taken.
 ✗ if (result != 0) {
236 ✗ CloseFd(fd);
237 }
238 ✗ m_hip.Check(result, "hipImportExternalSemaphore");
239 ✗ return imported;
240 }
241
242 const hip::Api& m_hip;
243 int m_device = 0;
244 hip::hipStream_t m_stream = nullptr;
245 hip::hipExternalSemaphore_t m_timeline = nullptr;
246 hip::hipExternalSemaphore_t m_follow = nullptr;
247 };
248 #endif
249
250 class HostInterop : public StreamInterop {
251 public:
252 ✗ HostInterop(VulkanDevice& device, VkSemaphore timeline, VkSemaphore follow)
253
0/1
✗ Branch 3 not taken.
 ✗ : m_device(device), m_timeline(timeline), m_follow(follow) {}
254
255 ✗ void Follow(void* stream, uint64_t value) override {
256 ✗ void* event = Record(stream);
257 ✗ {
258 ✗ const std::lock_guard lock(m_mutex);
259
0/1
✗ Branch 1 not taken.
 ✗ m_pending.push_back({event, value});
260
0/2
✗ Branch 0 not taken.
✗ Branch 1 not taken.
 ✗ if (!m_worker.joinable()) {
261
0/1
✗ Branch 2 not taken.
 ✗ m_worker = std::thread([this] { Run(); });
262 }
263 ✗ }
264 ✗ m_wake.notify_one();
265 ✗ }
266
267 ✗ void Wait(void* stream, uint64_t value) override {
268 ✗ auto* waiting = new Waiting{this, value};
269 ✗ {
270 ✗ const std::lock_guard lock(m_mutex);
271 ✗ ++m_waiting;
272 ✗ }
273 ✗ try {
274
0/1
✗ Branch 1 not taken.
 ✗ Launch(stream, WaitTimeline, waiting);
275 ✗ } catch (...) {
276
0/2
✗ Branch 0 not taken.
✗ Branch 1 not taken.
 ✗ delete waiting;
277
0/1
✗ Branch 1 not taken.
 ✗ Done();
278 ✗ throw;
279 ✗ }
280 ✗ }
281
282 protected:
283 virtual void* Record(void* stream) = 0;
284 virtual void Synchronize(void* event) = 0;
285 virtual void Launch(void* stream, void (*function)(void*), void* data) = 0;
286
287 ✗ void Stop() {
288 ✗ {
289 ✗ std::unique_lock lock(m_mutex);
290 ✗ m_stopping = true;
291
0/4
✗ Branch 0 not taken.
✗ Branch 1 not taken.
✗ Branch 2 not taken.
✗ Branch 3 not taken.
 ✗ m_idle.wait(lock, [this] { return m_waiting == 0; });
292 ✗ }
293 ✗ m_wake.notify_one();
294
0/2
✗ Branch 0 not taken.
✗ Branch 1 not taken.
 ✗ if (m_worker.joinable()) {
295 ✗ m_worker.join();
296 }
297 ✗ }
298
299 private:
300 struct Pending {
301 void* event = nullptr;
302 uint64_t value = 0;
303 };
304
305 struct Waiting {
306 HostInterop* interop = nullptr;
307 uint64_t value = 0;
308 };
309
310 ✗ static void WaitTimeline(void* data) {
311
0/1
✗ Branch 1 not taken.
 ✗ const std::unique_ptr<Waiting> waiting(static_cast<Waiting*>(data));
312
0/1
✗ Branch 1 not taken.
 ✗ HostInterop& interop = *waiting->interop;
313 ✗ VkSemaphoreWaitInfo wait{VK_STRUCTURE_TYPE_SEMAPHORE_WAIT_INFO};
314 ✗ wait.semaphoreCount = 1;
315 ✗ wait.pSemaphores = &interop.m_timeline;
316 ✗ wait.pValues = &waiting->value;
317 ✗ const VkResult result = interop.m_device.Functions().vkWaitSemaphores(
318
0/1
✗ Branch 1 not taken.
 ✗ interop.m_device.Handle(), &wait, UINT64_MAX);
319
0/2
✗ Branch 0 not taken.
✗ Branch 1 not taken.
 ✗ if (result != VK_SUCCESS) {
320 ✗ Log(LogLevel::Warning, "vkWaitSemaphores failed with {} while a stream waited",
321
0/1
✗ Branch 1 not taken.
 ✗ int(result));
322 }
323
0/1
✗ Branch 1 not taken.
 ✗ interop.Done();
324 ✗ }
325
326 ✗ void Done() {
327 ✗ const std::lock_guard lock(m_mutex);
328 ✗ --m_waiting;
329 ✗ m_idle.notify_all();
330 ✗ }
331
332 ✗ void Run() {
333 ✗ std::unique_lock lock(m_mutex);
334 ✗ while (true) {
335
0/5
✗ Branch 0 not taken.
✗ Branch 1 not taken.
✗ Branch 2 not taken.
✗ Branch 3 not taken.
✗ Branch 5 not taken.
 ✗ m_wake.wait(lock, [this] { return m_stopping || !m_pending.empty(); });
336
0/2
✗ Branch 0 not taken.
✗ Branch 1 not taken.
 ✗ if (m_pending.empty()) {
337
0/2
✗ Branch 0 not taken.
✗ Branch 1 not taken.
 ✗ return;
338 }
339 ✗ const Pending next = m_pending.front();
340 ✗ m_pending.pop_front();
341
0/1
✗ Branch 1 not taken.
 ✗ lock.unlock();
342 ✗ try {
343
0/1
✗ Branch 1 not taken.
 ✗ Synchronize(next.event);
344
0/2
✗ Branch 0 not taken.
✗ Branch 1 not taken.
 ✗ } catch (const Error& error) {
345
0/1
✗ Branch 2 not taken.
 ✗ Log(LogLevel::Warning, "A followed stream failed: {}", error.what());
346 ✗ }
347 ✗ VkSemaphoreSignalInfo signal{VK_STRUCTURE_TYPE_SEMAPHORE_SIGNAL_INFO};
348 ✗ signal.semaphore = m_follow;
349 ✗ signal.value = next.value;
350 ✗ const VkResult result =
351
0/1
✗ Branch 1 not taken.
 ✗ m_device.Functions().vkSignalSemaphore(m_device.Handle(), &signal);
352
0/2
✗ Branch 0 not taken.
✗ Branch 1 not taken.
 ✗ if (result != VK_SUCCESS) {
353
0/3
✗ Branch 1 not taken.
✗ Branch 3 not taken.
✗ Branch 4 not taken.
 ✗ Log(LogLevel::Warning, "vkSignalSemaphore failed with {}", int(result));
354 }
355
0/1
✗ Branch 1 not taken.
 ✗ lock.lock();
356 ✗ }
357 ✗ }
358
359 VulkanDevice& m_device;
360 VkSemaphore m_timeline;
361 VkSemaphore m_follow;
362 std::mutex m_mutex;
363 std::condition_variable m_wake;
364 std::condition_variable m_idle;
365 std::deque<Pending> m_pending;
366 uint32_t m_waiting = 0;
367 bool m_stopping = false;
368 std::thread m_worker;
369 };
370
371 class CudaHostInterop : public HostInterop {
372 public:
373 ✗ CudaHostInterop(VulkanDevice& device, VkSemaphore timeline, VkSemaphore follow)
374
0/1
✗ Branch 2 not taken.
 ✗ : HostInterop(device, timeline, follow), m_cuda(cuda::GetApi()) {
375
0/1
✗ Branch 1 not taken.
 ✗ m_device = FindCudaDevice(m_cuda, device).device;
376
0/2
✗ Branch 1 not taken.
✗ Branch 4 not taken.
 ✗ m_cuda.Check(m_cuda.cuDevicePrimaryCtxRetain(&m_context, m_device),
377 "cuDevicePrimaryCtxRetain");
378 ✗ }
379
380 ✗ ~CudaHostInterop() override {
381 ✗ Stop();
382 ✗ m_cuda.Warn(m_cuda.cuDevicePrimaryCtxRelease(m_device), "cuDevicePrimaryCtxRelease");
383 ✗ }
384
385 protected:
386 ✗ void* Record(void* stream) override {
387 ✗ const CudaContext scope(m_cuda, m_context);
388
0/1
✗ Branch 1 not taken.
 ✗ const cuda::CUevent event = cuda::Traits::CreateEvent(m_cuda, false);
389 ✗ try {
390
0/2
✗ Branch 1 not taken.
✗ Branch 4 not taken.
 ✗ cuda::Traits::Record(m_cuda, event, cuda::Traits::ApplicationStream(stream), false);
391 ✗ } catch (...) {
392
0/1
✗ Branch 1 not taken.
 ✗ cuda::Traits::DestroyEvent(m_cuda, event);
393 ✗ throw;
394 ✗ }
395 ✗ return event;
396 ✗ }
397
398 ✗ void Synchronize(void* event) override {
399 ✗ const CudaContext scope(m_cuda, m_context);
400 ✗ const auto recorded = static_cast<cuda::CUevent>(event);
401
0/1
✗ Branch 1 not taken.
 ✗ cuda::Traits::Synchronize(m_cuda, recorded);
402
0/1
✗ Branch 1 not taken.
 ✗ cuda::Traits::DestroyEvent(m_cuda, recorded);
403 ✗ }
404
405 ✗ void Launch(void* stream, void (*function)(void*), void* data) override {
406 ✗ const CudaContext scope(m_cuda, m_context);
407
0/2
✗ Branch 1 not taken.
✗ Branch 4 not taken.
 ✗ m_cuda.Check(
408
0/1
✗ Branch 1 not taken.
 ✗ m_cuda.cuLaunchHostFunc(cuda::Traits::ApplicationStream(stream), function, data),
409 "cuLaunchHostFunc");
410 ✗ }
411
412 private:
413 const cuda::Api& m_cuda;
414 cuda::CUdevice m_device = 0;
415 cuda::CUcontext m_context = nullptr;
416 };
417
418 class HipHostInterop : public HostInterop {
419 public:
420 ✗ HipHostInterop(VulkanDevice& device, VkSemaphore timeline, VkSemaphore follow)
421 ✗ : HostInterop(device, timeline, follow), m_hip(hip::GetApi()),
422
0/2
✗ Branch 2 not taken.
✗ Branch 5 not taken.
 ✗ m_device(FindHipDevice(m_hip, device)) {}
423
424 ✗ ~HipHostInterop() override {
425 ✗ Stop();
426 ✗ }
427
428 protected:
429 ✗ void* Record(void* stream) override {
430 ✗ Enter();
431 ✗ const hip::hipEvent_t event = hip::Traits::CreateEvent(m_hip, false);
432 ✗ try {
433
0/2
✗ Branch 1 not taken.
✗ Branch 4 not taken.
 ✗ hip::Traits::Record(m_hip, event, hip::Traits::ApplicationStream(stream), false);
434 ✗ } catch (...) {
435
0/1
✗ Branch 1 not taken.
 ✗ hip::Traits::DestroyEvent(m_hip, event);
436 ✗ throw;
437 ✗ }
438 ✗ return event;
439 }
440
441 ✗ void Synchronize(void* event) override {
442 ✗ Enter();
443 ✗ const auto recorded = static_cast<hip::hipEvent_t>(event);
444 ✗ hip::Traits::Synchronize(m_hip, recorded);
445 ✗ hip::Traits::DestroyEvent(m_hip, recorded);
446 ✗ }
447
448 ✗ void Launch(void* stream, void (*function)(void*), void* data) override {
449 ✗ Enter();
450 ✗ m_hip.Check(m_hip.hipLaunchHostFunc(hip::Traits::ApplicationStream(stream), function, data),
451 "hipLaunchHostFunc");
452 ✗ }
453
454 private:
455 ✗ void Enter() const {
456 ✗ m_hip.Check(m_hip.hipSetDevice(m_device), "hipSetDevice");
457 ✗ }
458
459 const hip::Api& m_hip;
460 int m_device = 0;
461 };
462
463 class CudaMemoryInterop : public MemoryInterop {
464 public:
465
0/1
✗ Branch 1 not taken.
 ✗ explicit CudaMemoryInterop(VulkanDevice& device) : m_cuda(cuda::GetApi()) {
466
0/1
✗ Branch 1 not taken.
 ✗ const CudaDevice found = FindCudaDevice(m_cuda, device);
467 ✗ m_device = found.device;
468 ✗ m_ordinal = found.ordinal;
469
0/2
✗ Branch 1 not taken.
✗ Branch 4 not taken.
 ✗ m_cuda.Check(m_cuda.cuDevicePrimaryCtxRetain(&m_context, m_device),
470 "cuDevicePrimaryCtxRetain");
471
0/1
✗ Branch 1 not taken.
 ✗ const CudaContext scope(m_cuda, m_context);
472
0/1
✗ Branch 1 not taken.
 ✗ m_stream = cuda::Traits::CreateStream(m_cuda);
473 ✗ }
474
475 ✗ ~CudaMemoryInterop() override {
476 ✗ {
477 ✗ const CudaContext scope(m_cuda, m_context);
478
0/2
✗ Branch 0 not taken.
✗ Branch 1 not taken.
 ✗ if (m_stream) {
479 ✗ cuda::Traits::DestroyStream(m_cuda, m_stream);
480 }
481 ✗ }
482 ✗ m_cuda.Warn(m_cuda.cuDevicePrimaryCtxRelease(m_device), "cuDevicePrimaryCtxRelease");
483 ✗ }
484
485 ✗ DLDevice Device() const override {
486 ✗ return {kDLCUDA, m_ordinal};
487 }
488
489 ✗ ImportedMemory Import(const ExportedMemory& memory) override {
490 ✗ const CudaContext scope(m_cuda, m_context);
491 ✗ cuda::CUDA_EXTERNAL_MEMORY_HANDLE_DESC desc{};
492 ✗ desc.type = cuda::CU_EXTERNAL_MEMORY_HANDLE_TYPE_OPAQUE_FD;
493 ✗ desc.handle.fd = memory.fd;
494 ✗ desc.size = memory.memoryBytes;
495 ✗ desc.flags = cuda::CUDA_EXTERNAL_MEMORY_DEDICATED;
496 ✗ cuda::CUexternalMemory external = nullptr;
497
0/1
✗ Branch 1 not taken.
 ✗ const cuda::CUresult imported = m_cuda.cuImportExternalMemory(&external, &desc);
498
0/2
✗ Branch 0 not taken.
✗ Branch 1 not taken.
 ✗ if (imported != cuda::CUDA_SUCCESS) {
499
0/1
✗ Branch 1 not taken.
 ✗ CloseFd(memory.fd);
500 }
501
0/1
✗ Branch 1 not taken.
 ✗ m_cuda.Check(imported, "cuImportExternalMemory");
502
503 ✗ cuda::CUDA_EXTERNAL_MEMORY_BUFFER_DESC buffer{};
504 ✗ buffer.offset = memory.offset;
505 ✗ buffer.size = memory.bytes;
506 ✗ cuda::CUdeviceptr address = 0;
507 ✗ const cuda::CUresult mapped =
508
0/1
✗ Branch 1 not taken.
 ✗ m_cuda.cuExternalMemoryGetMappedBuffer(&address, external, &buffer);
509
0/2
✗ Branch 0 not taken.
✗ Branch 1 not taken.
 ✗ if (mapped != cuda::CUDA_SUCCESS) {
510
0/2
✗ Branch 1 not taken.
✗ Branch 4 not taken.
 ✗ m_cuda.Warn(m_cuda.cuDestroyExternalMemory(external), "cuDestroyExternalMemory");
511 }
512
0/1
✗ Branch 1 not taken.
 ✗ m_cuda.Check(mapped, "cuExternalMemoryGetMappedBuffer");
513 ✗ return {external, address};
514 ✗ }
515
516 ✗ void Release(ImportedMemory& memory) noexcept override {
517
0/2
✗ Branch 0 not taken.
✗ Branch 1 not taken.
 ✗ if (memory.external == nullptr) {
518 return;
519 }
520 ✗ try {
521
0/1
✗ Branch 1 not taken.
 ✗ const CudaContext scope(m_cuda, m_context);
522
0/2
✗ Branch 1 not taken.
✗ Branch 4 not taken.
 ✗ m_cuda.Warn(m_cuda.cuMemFree(memory.address), "cuMemFree");
523
0/1
✗ Branch 1 not taken.
 ✗ m_cuda.Warn(m_cuda.cuDestroyExternalMemory(
524
0/1
✗ Branch 1 not taken.
 ✗ static_cast<cuda::CUexternalMemory>(memory.external)),
525 "cuDestroyExternalMemory");
526 ✗ } catch (...) {
527 ✗ }
528 ✗ memory = {};
529 }
530
531 ✗ void* Stream() const override {
532 ✗ return m_stream;
533 }
534
535 ✗ void Copy(uint64_t dst, uint64_t src, uint64_t bytes, void* stream) override {
536 ✗ const CudaContext scope(m_cuda, m_context);
537
0/3
✗ Branch 0 not taken.
✗ Branch 1 not taken.
✗ Branch 3 not taken.
 ✗ const cuda::CUstream on = stream ? cuda::Traits::ApplicationStream(stream) : m_stream;
538
0/2
✗ Branch 1 not taken.
✗ Branch 4 not taken.
 ✗ m_cuda.Check(m_cuda.cuMemcpyAsync(dst, src, bytes, on), "cuMemcpyAsync");
539
0/2
✗ Branch 0 not taken.
✗ Branch 1 not taken.
 ✗ if (stream == nullptr) {
540
0/2
✗ Branch 1 not taken.
✗ Branch 4 not taken.
 ✗ m_cuda.Check(m_cuda.cuStreamSynchronize(m_stream), "cuStreamSynchronize");
541 }
542 ✗ }
543
544 private:
545 const cuda::Api& m_cuda;
546 cuda::CUdevice m_device = 0;
547 int32_t m_ordinal = 0;
548 cuda::CUcontext m_context = nullptr;
549 cuda::CUstream m_stream = nullptr;
550 };
551
552 class HipMemoryInterop : public MemoryInterop {
553 public:
554
0/1
✗ Branch 1 not taken.
 ✗ explicit HipMemoryInterop(VulkanDevice& device) : m_hip(hip::GetApi()) {
555
0/1
✗ Branch 1 not taken.
 ✗ m_device = FindHipDevice(m_hip, device);
556
0/1
✗ Branch 1 not taken.
 ✗ Enter();
557
0/1
✗ Branch 1 not taken.
 ✗ m_stream = hip::Traits::CreateStream(m_hip);
558 ✗ }
559
560 ✗ ~HipMemoryInterop() override {
561 ✗ m_hip.Warn(m_hip.hipSetDevice(m_device), "hipSetDevice");
562
0/2
✗ Branch 0 not taken.
✗ Branch 1 not taken.
 ✗ if (m_stream) {
563 ✗ hip::Traits::DestroyStream(m_hip, m_stream);
564 }
565 ✗ }
566
567 ✗ DLDevice Device() const override {
568 ✗ return {kDLROCM, m_device};
569 }
570
571 ✗ ImportedMemory Import(const ExportedMemory& memory) override {
572 ✗ Enter();
573 ✗ hip::hipExternalMemoryHandleDesc desc{};
574 ✗ desc.type = hip::hipExternalMemoryHandleTypeOpaqueFd;
575 ✗ desc.handle.fd = memory.fd;
576 ✗ desc.size = memory.memoryBytes;
577 ✗ desc.flags = hip::hipExternalMemoryDedicated;
578 ✗ hip::hipExternalMemory_t external = nullptr;
579 ✗ const hip::hipError_t imported = m_hip.hipImportExternalMemory(&external, &desc);
580
0/2
✗ Branch 0 not taken.
✗ Branch 1 not taken.
 ✗ if (imported != hip::hipSuccess) {
581 ✗ CloseFd(memory.fd);
582 }
583 ✗ m_hip.Check(imported, "hipImportExternalMemory");
584
585 ✗ hip::hipExternalMemoryBufferDesc buffer{};
586 ✗ buffer.offset = memory.offset;
587 ✗ buffer.size = memory.bytes;
588 ✗ void* address = nullptr;
589 ✗ const hip::hipError_t mapped =
590 ✗ m_hip.hipExternalMemoryGetMappedBuffer(&address, external, &buffer);
591
0/2
✗ Branch 0 not taken.
✗ Branch 1 not taken.
 ✗ if (mapped != hip::hipSuccess) {
592 ✗ m_hip.Warn(m_hip.hipDestroyExternalMemory(external), "hipDestroyExternalMemory");
593 }
594 ✗ m_hip.Check(mapped, "hipExternalMemoryGetMappedBuffer");
595 ✗ return {external, uint64_t(address)};
596 }
597
598 ✗ void Release(ImportedMemory& memory) noexcept override {
599
0/2
✗ Branch 0 not taken.
✗ Branch 1 not taken.
 ✗ if (memory.external == nullptr) {
600 return;
601 }
602 ✗ try {
603
0/1
✗ Branch 1 not taken.
 ✗ Enter();
604
0/2
✗ Branch 1 not taken.
✗ Branch 4 not taken.
 ✗ m_hip.Warn(m_hip.hipFree(reinterpret_cast<void*>(memory.address)), "hipFree");
605
0/2
✗ Branch 1 not taken.
✗ Branch 4 not taken.
 ✗ m_hip.Warn(m_hip.hipDestroyExternalMemory(memory.external), "hipDestroyExternalMemory");
606 ✗ } catch (...) {
607 ✗ }
608 ✗ memory = {};
609 }
610
611 ✗ void* Stream() const override {
612 ✗ return m_stream;
613 }
614
615 ✗ void Copy(uint64_t dst, uint64_t src, uint64_t bytes, void* stream) override {
616 ✗ Enter();
617
0/2
✗ Branch 0 not taken.
✗ Branch 1 not taken.
 ✗ const hip::hipStream_t on = stream ? hip::Traits::ApplicationStream(stream) : m_stream;
618 ✗ m_hip.Check(m_hip.hipMemcpyAsync(reinterpret_cast<void*>(dst),
619 reinterpret_cast<const void*>(src), bytes,
620 hip::hipMemcpyDefault, on),
621 "hipMemcpyAsync");
622
0/2
✗ Branch 0 not taken.
✗ Branch 1 not taken.
 ✗ if (stream == nullptr) {
623 ✗ m_hip.Check(m_hip.hipStreamSynchronize(m_stream), "hipStreamSynchronize");
624 }
625 ✗ }
626
627 private:
628 ✗ void Enter() const {
629 ✗ m_hip.Check(m_hip.hipSetDevice(m_device), "hipSetDevice");
630 ✗ }
631
632 const hip::Api& m_hip;
633 int m_device = 0;
634 hip::hipStream_t m_stream = nullptr;
635 };
636
637 }
638
639 ✗ std::unique_ptr<StreamInterop> CreateStreamInterop(VulkanDevice& device, VkSemaphore timeline,
640 VkSemaphore follow) {
641
0/4
✗ Branch 0 not taken.
✗ Branch 1 not taken.
✗ Branch 3 not taken.
✗ Branch 4 not taken.
 ✗ if (device.VendorId() == kVendorNvidia && cuda::Available()) {
642 #if !defined(_WIN32)
643
0/2
✗ Branch 0 not taken.
✗ Branch 1 not taken.
 ✗ if (device.SharesSemaphores()) {
644 ✗ return std::make_unique<CudaInterop>(device, timeline, follow);
645 }
646 #endif
647 ✗ return std::make_unique<CudaHostInterop>(device, timeline, follow);
648 }
649
0/4
✗ Branch 0 not taken.
✗ Branch 1 not taken.
✗ Branch 3 not taken.
✗ Branch 4 not taken.
 ✗ if (device.VendorId() == kVendorAmd && hip::Available()) {
650 #if !defined(_WIN32)
651
0/2
✗ Branch 0 not taken.
✗ Branch 1 not taken.
 ✗ if (device.SharesSemaphores()) {
652 ✗ return std::make_unique<HipInterop>(device, timeline, follow);
653 }
654 #endif
655 ✗ return std::make_unique<HipHostInterop>(device, timeline, follow);
656 }
657 ✗ throw UnsupportedError("{} shares no streams with CUDA or HIP: its GPU is of neither here",
658
0/1
✗ Branch 2 not taken.
 ✗ device.Adapter().alias);
659 }
660
661 13 std::unique_ptr<MemoryInterop> CreateMemoryInterop(VulkanDevice& device) {
662
1/2
✗ Branch 0 not taken.
✓ Branch 1 taken 13 times.
 13 if (!device.SharesExportedMemory()) {
663
0/1
✗ Branch 2 not taken.
 ✗ throw UnsupportedError("{} exports no memory", device.Adapter().alias);
664 }
665
1/4
✗ Branch 0 not taken.
✓ Branch 1 taken 13 times.
✗ Branch 3 not taken.
✗ Branch 4 not taken.
 13 if (device.VendorId() == kVendorNvidia && cuda::Available()) {
666 ✗ return std::make_unique<CudaMemoryInterop>(device);
667 }
668
1/4
✗ Branch 0 not taken.
✓ Branch 1 taken 13 times.
✗ Branch 3 not taken.
✗ Branch 4 not taken.
 13 if (device.VendorId() == kVendorAmd && hip::Available()) {
669 ✗ return std::make_unique<HipMemoryInterop>(device);
670 }
671 13 throw UnsupportedError("{} shares no memory with CUDA or HIP: its GPU is of neither here",
672
1/1
✓ Branch 2 taken 13 times.
 13 device.Adapter().alias);
673 }
674
675 }
676