GCC Code Coverage Report

Directory: avs_core/
Coverage: low: ≥ 0% medium: ≥ 75.0% high: ≥ 90.0%
Coverage Exec / Excl / Total
Lines: 19.5% 49 / 0 / 251
Functions: 25.7% 9 / 0 / 35
Branches: 8.2% 24 / 0 / 294
core/DeviceManager.cpp
Line Branch Exec Source
1
2 #include "DeviceManager.h"
3 #include "internal.h"
4 #include "InternalEnvironment.h"
5 #include <avs/minmax.h>
6 #include <deque>
7 #include <map>
8 #include <mutex>
9 #include <sstream>
10 #include "LruCache.h"
11 #include "ThreadPool.h"
12 #include "AVSMap.h"
13 #include "parser/scriptparser.h"
14
15 #ifndef MINGW_HAS_SECURE_API
16 #define sprintf_s sprintf
17 #endif
18
19 #define ENABLE_CUDA_COMPUTE_STREAM 0
20
21 #ifdef ENABLE_CUDA
22
23 #include <cuda_runtime_api.h>
24
25 #endif // #ifdef ENABLE_CUDA
26
27 #define CUDA_CHECK(call) \
28 do { \
29 cudaError_t err__ = call; \
30 if (err__ != cudaSuccess) { \
31 env->ThrowError("[CUDA Error] %d: %s @%d", err__, cudaGetErrorString(err__), __LINE__); \
32 } \
33 } while (0)
34
35 19 int GetDeviceTypes(const PClip& child)
36 {
37
1/2
✗ Branch 4 → 5 not taken.
✓ Branch 4 → 6 taken 19 times.
 19 if (child->GetVersion() < 5) {
38 ✗ return DEV_TYPE_CPU;
39 }
40 19 int deviceflags = child->SetCacheHints(CACHE_GET_DEV_TYPE, 0);
41
1/2
✓ Branch 8 → 9 taken 19 times.
✗ Branch 8 → 10 not taken.
 19 if (deviceflags == 0) {
42 // if not implement CACHE_GET_DEVICE_TYPE, we assume CPU only filter.
43 19 deviceflags = DEV_TYPE_CPU;
44 }
45 19 return deviceflags;
46 }
47
48 7 int GetTargetDeviceTypes(const PClip& clip)
49 {
50
1/2
✗ Branch 4 → 5 not taken.
✓ Branch 4 → 6 taken 7 times.
 7 if (clip->GetVersion() < 5) {
51 ✗ return DEV_TYPE_CPU;
52 }
53 7 int deviceflags = clip->SetCacheHints(CACHE_GET_CHILD_DEV_TYPE, 0);
54
1/2
✓ Branch 8 → 9 taken 7 times.
✗ Branch 8 → 13 not taken.
 7 if (deviceflags == 0) {
55 7 deviceflags = clip->SetCacheHints(CACHE_GET_DEV_TYPE, 0);
56
1/2
✓ Branch 11 → 12 taken 7 times.
✗ Branch 11 → 13 not taken.
 7 if (deviceflags == 0) {
57 // if not implement CACHE_GET_DEVICE_TYPE, we assume CPU only filter.
58 7 deviceflags = DEV_TYPE_CPU;
59 }
60 }
61 7 return deviceflags;
62 }
63
64 ✗ std::string DeviceTypesString(int devicetypes)
65 {
66 ✗ std::vector<const char*> typesstr;
67 ✗ if (devicetypes & DEV_TYPE_CPU) {
68 ✗ typesstr.push_back("CPU");
69 }
70 ✗ if (devicetypes & DEV_TYPE_CUDA) {
71 ✗ typesstr.push_back("CUDA");
72 }
73 ✗ std::ostringstream oss;
74 ✗ for (int i = 0; i < (int)typesstr.size(); ++i) {
75 ✗ if (i > 0) oss << ",";
76 ✗ oss << typesstr[i];
77 }
78 ✗ return oss.str();
79 ✗ }
80
81 7 static void CheckDeviceTypes(const char* name, int devicetypes, const AVSValue& last, const AVSValue& arr, InternalEnvironment* env)
82 {
83
2/2
✓ Branch 29 → 3 taken 23 times.
✓ Branch 29 → 30 taken 7 times.
 30 for (int i = -1; i < arr.ArraySize(); ++i) {
84
2/2
✓ Branch 3 → 4 taken 16 times.
✓ Branch 3 → 6 taken 7 times.
 23 const AVSValue& val = (i == -1) ? last : arr[i];
85
2/2
✓ Branch 8 → 9 taken 7 times.
✓ Branch 8 → 21 taken 16 times.
 23 if (val.IsClip()) {
86
2/4
✓ Branch 9 → 10 taken 7 times.
✗ Branch 9 → 33 not taken.
✓ Branch 10 → 11 taken 7 times.
✗ Branch 10 → 31 not taken.
 7 int childtypes = GetDeviceTypes(val.AsClip());
87
1/2
✗ Branch 12 → 13 not taken.
✓ Branch 12 → 27 taken 7 times.
 7 if ((devicetypes & childtypes) == 0) {
88 ✗ std::string parentdevstr = DeviceTypesString(devicetypes);
89 ✗ std::string childdevstr = DeviceTypesString(childtypes);
90 // e.g.. Device unmatch: XYfilter[CPU] does not support [CUDA] frame
91 ✗ env->ThrowError(
92 "Device unmatch: %s[%s] does not support [%s] frame",
93 name, parentdevstr.c_str(), childdevstr.c_str());
94 ✗ }
95 }
96
1/2
✗ Branch 22 → 23 not taken.
✓ Branch 22 → 27 taken 16 times.
 16 else if (val.IsArray()) {
97 ✗ CheckDeviceTypes(name, devicetypes, AVSValue(), val, env);
98 }
99 }
100 7 }
101
102 7 void CheckChildDeviceTypes(const PClip& clip, const char* name, const AVSValue& last, const AVSValue& args, const char* const* argnames, InternalEnvironment* env)
103 {
104 7 int deviceflags = GetTargetDeviceTypes(clip);
105
2/2
✓ Branch 4 → 5 taken 6 times.
✓ Branch 4 → 6 taken 1 time.
 7 if (args.IsArray()) {
106 6 CheckDeviceTypes(name, deviceflags, last, args, env);
107 }
108 else {
109
2/4
✓ Branch 6 → 7 taken 1 time.
✗ Branch 6 → 13 not taken.
✓ Branch 7 → 8 taken 1 time.
✗ Branch 7 → 11 not taken.
 1 CheckDeviceTypes(name, deviceflags, last, AVSValue(&args, 1), env);
110 }
111 7 }
112
113 ✗ size_t GetFrameHead(const PVideoFrame& vf)
114 {
115 ✗ int head = vf->GetOffset();
116 ✗ if (vf->GetPitch(PLANAR_U)) {
117 ✗ head = min(head, vf->GetOffset(PLANAR_U));
118 }
119 ✗ if (vf->GetPitch(PLANAR_V)) {
120 ✗ head = min(head, vf->GetOffset(PLANAR_V));
121 }
122 ✗ if (vf->GetPitch(PLANAR_A)) {
123 ✗ head = min(head, vf->GetOffset(PLANAR_A));
124 }
125 ✗ return head;
126 }
127
128 ✗ size_t GetFrameTail(const PVideoFrame& vf)
129 {
130 ✗ int tail = vf->GetOffset() + vf->GetPitch() * vf->GetHeight();
131 ✗ if (vf->GetPitch(PLANAR_U)) {
132 ✗ tail = max(tail, vf->GetOffset(PLANAR_U) + vf->GetPitch(PLANAR_U) * vf->GetHeight(PLANAR_U));
133 }
134 ✗ if (vf->GetPitch(PLANAR_V)) {
135 ✗ tail = max(tail, vf->GetOffset(PLANAR_V) + vf->GetPitch(PLANAR_V) * vf->GetHeight(PLANAR_V));
136 }
137 ✗ if (vf->GetPitch(PLANAR_A)) {
138 ✗ tail = max(tail, vf->GetOffset(PLANAR_A) + vf->GetPitch(PLANAR_A) * vf->GetHeight(PLANAR_A));
139 }
140 ✗ return min(tail, vf->GetFrameBuffer()->GetDataSize() + (int)GetFrameHead(vf));
141 }
142
143 class CPUDevice : public Device {
144 public:
145 453 CPUDevice(InternalEnvironment* env)
146 453 : Device(DEV_TYPE_CPU, 0, 0, env)
147 453 { }
148
149 ✗ virtual int SetMemoryMax(int mem)
150 {
151 // memory_max for CPU device is not implemented here.
152 ✗ env->ThrowError("Not implemented ...");
153 ✗ return 0;
154 }
155
156 891 virtual BYTE* Allocate(size_t size, int margin)
157 {
158 891 size += margin;
159 // Allocate a bit more, so we can safely use a 64 byte (512 bits AVX512) load even from
160 // the last pixel of VideoFrameBuffer. Note: NaN contamination over valid scanlines are still a concern
161 // when using 32-bit float data.
162 // Pre 3.7.6 used only 16 bytes (SSE2)
163 891 constexpr size_t SIMD_OVERREAD_EXTRA_SAFETY = 64;
164 #ifdef _DEBUG
165 BYTE* data = new BYTE[size + SIMD_OVERREAD_EXTRA_SAFETY];
166 // fill extra space with recognizable pattern
167 for (auto i = 0; i < SIMD_OVERREAD_EXTRA_SAFETY / 8; ++i) {
168 ((uint64_t*)(data + size))[i] = 0xDEADBEEFDEADBEEF;
169 }
170 // fill allocated memory with recognizable pattern
171 // AllOCA7ED, uneven 5-byte pattern to catch overruns
172 static const BYTE filler[] = { 0x0A, 0x11, 0x0C, 0xA7, 0xED };
173 BYTE* pByte = data;
174 BYTE* q = pByte + size / 5 * 5;
175 for (; pByte < q; pByte += 5)
176 {
177 pByte[0] = filler[0];
178 pByte[1] = filler[1];
179 pByte[2] = filler[2];
180 pByte[3] = filler[3];
181 pByte[4] = filler[4];
182 }
183 return data;
184 #else
185 891 return new BYTE[size + SIMD_OVERREAD_EXTRA_SAFETY];
186 #endif
187 }
188
189 891 virtual void Free(BYTE* ptr)
190 {
191
1/2
✓ Branch 2 → 3 taken 891 times.
✗ Branch 2 → 5 not taken.
 891 if (ptr != nullptr) {
192
1/2
✓ Branch 3 → 4 taken 891 times.
✗ Branch 3 → 5 not taken.
 891 delete[] ptr;
193 }
194 891 }
195
196 ✗ virtual const char* GetName() const { return "CPU"; }
197
198 ✗ virtual void AddCompleteCallback(DeviceCompleteCallbackData cbdata)
199 {
200 // no need to delay, call immediately
201 ✗ cbdata.cb(cbdata.user_data);
202 ✗ }
203
204 ✗ virtual std::unique_ptr<std::vector<DeviceCompleteCallbackData>> GetAndClearCallbacks()
205 {
206 ✗ return nullptr;
207 }
208
209 ✗ virtual void SetActiveToCurrentThread(InternalEnvironment* env)
210 {
211 // do nothing
212 ✗ }
213
214 ✗ virtual void* GetComputeStream()
215 {
216 ✗ return nullptr;
217 }
218
219 ✗ virtual void SetDeviceOpt(DeviceOpt opt, int val, InternalEnvironment* env)
220 {
221 // do nothing
222 ✗ }
223
224 ✗ virtual void GetAlignmentRequirement(int* memoryAlignment, int* pitchAlignment)
225 {
226 ✗ *memoryAlignment = FRAME_ALIGN;
227 ✗ *pitchAlignment = FRAME_ALIGN;
228 ✗ }
229 };
230
231 #ifdef ENABLE_CUDA
232 class CUDACPUDevice : public CPUDevice {
233 int num_cuda_devices;
234 bool allocated;
235 bool enable_pinned;
236 public:
237 CUDACPUDevice(InternalEnvironment* env, int num_cuda_devices)
238 : CPUDevice(env)
239 , num_cuda_devices(num_cuda_devices)
240 , allocated()
241 , enable_pinned()
242 { }
243
244 virtual BYTE* Allocate(size_t size, int margin)
245 {
246 allocated = true;
247 if (!enable_pinned) {
248 return CPUDevice::Allocate(size, margin);
249 }
250 //unsigned int flags = (prop.canMapHostMemory && prop.unifiedAddressing)
251 // ? cudaHostAllocMapped : cudaHostAllocDefault;
252 // Without portable flag, allocated memory is associated with one cuda context,
253 // i.e. one GPU. portable flag is required to use this memory with all GPUs.
254 unsigned int flags = (num_cuda_devices > 1)
255 ? cudaHostAllocPortable : cudaHostAllocDefault;
256 BYTE* data = nullptr;
257 size += margin;
258
259 // FIXME check: is extra safety needed for pinned memory like CPUDevice::Allocate
260 // adds extra 64 bytes for SIMD overread safety instead of 16 bytes?
261
262 #ifdef _DEBUG
263 CUDA_CHECK(cudaHostAlloc((void**)&data, size + 16, flags));
264 int *pInt = (int *)(data + size);
265 pInt[0] = 0xDEADBEEF;
266 pInt[1] = 0xDEADBEEF;
267 pInt[2] = 0xDEADBEEF;
268 pInt[3] = 0xDEADBEEF;
269
270 static const BYTE filler[] = { 0x0A, 0x11, 0x0C, 0xA7, 0xED };
271 BYTE* pByte = data;
272 BYTE* q = pByte + size / 5 * 5;
273 for (; pByte < q; pByte += 5)
274 {
275 pByte[0] = filler[0];
276 pByte[1] = filler[1];
277 pByte[2] = filler[2];
278 pByte[3] = filler[3];
279 pByte[4] = filler[4];
280 }
281 #else
282 CUDA_CHECK(cudaHostAlloc((void**)&data, size + 16, flags));
283 #endif
284 return data;
285 }
286
287 virtual void Free(BYTE* ptr)
288 {
289 if (!enable_pinned) {
290 CPUDevice::Free(ptr);
291 }
292 else {
293 if (ptr != nullptr) {
294 CUDA_CHECK(cudaFreeHost(ptr));
295 }
296 }
297 }
298
299 virtual const char* GetName() const { return "CPU(CUDAAware)"; }
300
301 void SetDeviceOpt(DeviceOpt opt, int val, InternalEnvironment* env)
302 {
303 if (opt == DEV_CUDA_PINNED_HOST) {
304 if (!enable_pinned) {
305 if (allocated) {
306 env->ThrowError("SetDeviceOpt: Allocation mode change must be done before any frame allocation.");
307 }
308 enable_pinned = true;
309 }
310 }
311 if (opt == DEV_FREE_THRESHOLD) {
312 free_thresh = val;
313 }
314 }
315
316 virtual void GetAlignmentRequirement(int* memoryAlignment, int* pitchAlignment)
317 {
318 *memoryAlignment = FRAME_ALIGN;
319 *pitchAlignment = FRAME_ALIGN;
320 }
321 };
322 #endif
323
324 #ifdef ENABLE_CUDA
325 class CUDADevice : public Device {
326 class ScopedCUDADevice
327 {
328 int old_device;
329 int tgt_device;
330 public:
331 ScopedCUDADevice(int device_index, IScriptEnvironment* env)
332 : tgt_device(device_index)
333 {
334 CUDA_CHECK(cudaGetDevice(&old_device));
335 if (tgt_device != old_device) {
336 CUDA_CHECK(cudaSetDevice(tgt_device));
337 }
338 }
339 ~ScopedCUDADevice()
340 {
341 if (tgt_device != old_device) {
342 cudaSetDevice(old_device);
343 }
344 }
345 };
346
347 char name[32];
348
349 cudaDeviceProp prop;
350
351 std::mutex mutex;
352 std::vector<DeviceCompleteCallbackData> callbacks;
353
354 #if ENABLE_CUDA_COMPUTE_STREAM
355 cudaStream_t computeStream;
356 cudaEvent_t computeEvent;
357 #endif
358
359 public:
360 CUDADevice(int id, int n, InternalEnvironment* env) :
361 Device(DEV_TYPE_CUDA, id, n, env)
362 {
363 sprintf_s(name, "CUDA %d", n);
364
365 ScopedCUDADevice d(device_index, env);
366
367 CUDA_CHECK(cudaGetDeviceProperties(&prop, device_index));
368
369 SetMemoryMax(768); // start with 768MB
370 #if ENABLE_CUDA_COMPUTE_STREAM
371 CUDA_CHECK(cudaStreamCreate(&computeStream));
372 CUDA_CHECK(cudaEventCreate(&computeEvent));
373 #endif
374 }
375
376 ~CUDADevice()
377 {
378 #if ENABLE_CUDA_COMPUTE_STREAM
379 cudaStreamDestroy(computeStream);
380 cudaEventDestroy(computeEvent);
381 #endif
382 }
383
384 virtual int SetMemoryMax(int mem)
385 {
386 if (mem > 0) {
387 uint64_t requested = mem * 1048576ull;
388 uint64_t mem_limit = prop.totalGlobalMem;
389 memory_max = clamp(requested, (uint64_t)(64 * 1024 * 1024ull), (uint64_t)(mem_limit - 128 * 1024 * 1024ull));
390 }
391 return (int)(memory_max / 1048576ull);
392 }
393
394 virtual BYTE* Allocate(size_t size, int margin)
395 {
396 BYTE* data = nullptr;
397 ScopedCUDADevice d(device_index, env);
398 CUDA_CHECK(cudaMalloc((void**)&data, size));
399 return data;
400 }
401
402 virtual void Free(BYTE* ptr)
403 {
404 if (ptr != NULL) {
405 ScopedCUDADevice d(device_index, env);
406 CUDA_CHECK(cudaFree(ptr));
407 }
408 }
409
410 virtual const char* GetName() const { return name; }
411
412 virtual void AddCompleteCallback(DeviceCompleteCallbackData cbdata)
413 {
414 std::lock_guard<std::mutex> lock(mutex);
415
416 callbacks.push_back(cbdata);
417 }
418
419 virtual std::unique_ptr<std::vector<DeviceCompleteCallbackData>> GetAndClearCallbacks()
420 {
421 std::lock_guard<std::mutex> lock(mutex);
422
423 if (callbacks.size() > 0) {
424 auto *ret = new std::vector<DeviceCompleteCallbackData>(std::move(callbacks));
425 callbacks.clear();
426 return std::unique_ptr<std::vector<DeviceCompleteCallbackData>>(ret);
427 }
428
429 return nullptr;
430 }
431
432 virtual void SetActiveToCurrentThread(InternalEnvironment* env)
433 {
434 CUDA_CHECK(cudaSetDevice(device_index));
435 }
436
437 virtual void* GetComputeStream() {
438 #if ENABLE_CUDA_COMPUTE_STREAM
439 return computeStream;
440 #else
441 return nullptr;
442 #endif
443 }
444
445 void MakeStreamWaitCompute(cudaStream_t stream, InternalEnvironment* env)
446 {
447 #if ENABLE_CUDA_COMPUTE_STREAM
448 std::lock_guard<std::mutex> lock(mutex);
449
450 CUDA_CHECK(cudaEventRecord(computeEvent, computeStream));
451 CUDA_CHECK(cudaStreamWaitEvent(stream, computeEvent, 0));
452 #endif
453 }
454
455 void SetDeviceOpt(DeviceOpt opt, int val, InternalEnvironment* env) {
456 if (opt == DEV_FREE_THRESHOLD) {
457 free_thresh = val;
458 }
459 }
460
461 virtual void GetAlignmentRequirement(int* memoryAlignment, int* pitchAlignment)
462 {
463 *memoryAlignment = (int)prop.textureAlignment;
464 *pitchAlignment = (int)prop.texturePitchAlignment;
465 }
466 };
467 #endif
468
469 453 DeviceManager::DeviceManager(InternalEnvironment* env) :
470 453 env(env)
471 {
472 // 0 is CPU device, start from 1 for other devices.
473 453 int next_device_id = 1;
474
475 #ifdef ENABLE_CUDA
476 int cuda_device_count = 0;
477 cudaError_t status = cudaGetDeviceCount(&cuda_device_count);
478 if (status == cudaSuccess) {
479 // _RPT0(0, "cudaGetDeviceCount = %d\r\n", cuda_device_count);
480 for (int i = 0; i < cuda_device_count; ++i) {
481 cudaDevices.emplace_back(new CUDADevice(next_device_id++, i, env));
482 }
483 }
484 else {
485 if (status == cudaErrorInitializationError) {
486 // We probably get this error because of current Nvidia driver version is lower than
487 // the minimum required version by the used CUDA SDK, or a Computing Capability is too low.
488 // Example: GTX460: Latest driver is 391.35, latest supporting CUDA SDK is 9.1.85
489 // Consequence: since SDK 9 is unsupported in VS2019, no GTX460 support
490 // As of Jan.2021 Avisynth+ is is using CUDA Toolkit version 11.2.
491 int version;
492 status = cudaRuntimeGetVersion(&version);
493 if(status == cudaSuccess)
494 _RPT1(0, "cudaGetDeviceCount: cudaErrorInitializationError!\r\nMaybe CUDA Runtime version (%d) does not support old drivers. \r\n", version);
495 else
496 _RPT0(0, "cudaGetDeviceCount: cudaErrorInitializationError! Runtime version request failed\r\n");
497 }
498 else {
499 _RPT1(0, "cudaGetDeviceCount failed (%d)\r\n", (int)status);
500 }
501 }
502 // do not modify CUDADevices after this since it causes pointer change
503
504 cpuDevice = std::unique_ptr<Device>((cuda_device_count > 0)
505 ? new CUDACPUDevice(env, cuda_device_count)
506 : new CPUDevice(env));
507 #else // ENABLE_CUDA
508
2/4
✓ Branch 4 → 5 taken 453 times.
✗ Branch 4 → 12 not taken.
✗ Branch 9 → 10 not taken.
✓ Branch 9 → 11 taken 453 times.
 453 cpuDevice = std::unique_ptr<CPUDevice>(new CPUDevice(env));
509 #endif // ENABLE_CUDA
510
511 453 numDevices = next_device_id;
512 453 }
513
514 961 Device* DeviceManager::GetDevice(AvsDeviceType device_type, int device_index) const
515 {
516
1/3
✓ Branch 2 → 3 taken 961 times.
✗ Branch 2 → 4 not taken.
✗ Branch 2 → 5 not taken.
 961 switch (device_type) {
517
518 961 case DEV_TYPE_CPU:
519 961 return cpuDevice.get();
520
521 ✗ case DEV_TYPE_CUDA:
522 #ifdef ENABLE_CUDA
523 if (device_index < 0) {
524 env->ThrowError("Invalid device index %d", device_index);
525 }
526 if (cudaDevices.size() == 0) {
527 env->ThrowError("No CUDA devices ...");
528 }
529 // wrap index
530 device_index %= (int)cudaDevices.size();
531 return cudaDevices[device_index].get();
532 #else
533 ✗ env->ThrowError("This Avisynth does not support memory type %d (CUDA)", device_type);
534 #endif
535
536 ✗ default:
537 ✗ env->ThrowError("Not supported memory type %d", device_type);
538 }
539 ✗ return nullptr;
540 }
541
542 ✗ int DeviceManager::GetNumDevices(AvsDeviceType device_type) const
543 {
544 ✗ switch (device_type) {
545
546 ✗ case DEV_TYPE_CPU:
547 ✗ return 1;
548
549 ✗ case DEV_TYPE_CUDA:
550 ✗ return (int)cudaDevices.size();
551
552 ✗ default:
553 ✗ env->ThrowError("Not supported memory type %d", device_type);
554 }
555 ✗ return 0;
556 }
557
558 ✗ void DeviceManager::SetDeviceOpt(DeviceOpt opt, int val, InternalEnvironment* env)
559 {
560 ✗ cpuDevice->SetDeviceOpt(opt, val, env);
561 #ifdef ENABLE_CUDA
562 for (auto& dev : cudaDevices) {
563 dev->SetDeviceOpt(opt, val, env);
564 }
565 #endif // #ifdef ENABLE_CUDA
566 ✗ }
567
568 ✗ DeviceSetter::DeviceSetter(InternalEnvironment* env, Device* upstreamDevice)
569 ✗ : env(env)
570 {
571 ✗ downstreamDevice = env->SetCurrentDevice(upstreamDevice);
572 ✗ if (downstreamDevice == nullptr) {
573 ✗ env->ThrowError("This thread is not created by AviSynth. It is not allowed to invoke script on this thread ...");
574 }
575 ✗ }
576
577 ✗ DeviceSetter::~DeviceSetter()
578 {
579 ✗ env->SetCurrentDevice(downstreamDevice);
580 ✗ }
581
582 class QueuePrefetcher
583 {
584 PClip child;
585 VideoInfo vi;
586
587 int prefetchFrames;
588 int numThreads;
589
590 ThreadPool* threadPool;
591 Device* device;
592
593 typedef LruCache<size_t, PVideoFrame> CacheType;
594
595 std::shared_ptr<CacheType> videoCache;
596
597 std::mutex mutex; // do not acceess to videoCache during locked by this mutex
598 std::deque<std::pair<size_t, CacheType::handle>> prefetchQueue;
599 int numWorkers;
600 std::exception_ptr workerException;
601 bool workerExceptionPresent;
602
603 static AVSValue ThreadWorker_(IScriptEnvironment2* env, void* data)
604 {
605 return static_cast<QueuePrefetcher*>(data)->ThreadWorker(
606 static_cast<InternalEnvironment*>(env));
607 }
608
609 AVSValue ThreadWorker(InternalEnvironment* env)
610 {
611 device->SetActiveToCurrentThread(env);
612
613 while (true) {
614 std::pair<size_t, CacheType::handle> work;
615 {
616 std::lock_guard<std::mutex> lock(mutex);
617 if (prefetchQueue.size() == 0) {
618 // there are no prefetch work
619 --numWorkers;
620 break;
621 }
622 work = prefetchQueue.front();
623 prefetchQueue.pop_front();
624 }
625
626 try
627 {
628 work.second.first->value = child->GetFrame((int)work.first, env);
629 videoCache->commit_value(&work.second);
630 }
631 catch (...)
632 {
633 {
634 std::lock_guard<std::mutex> lock(mutex);
635 workerException = std::current_exception();
636 workerExceptionPresent = true;
637 }
638 videoCache->rollback(&work.second);
639 }
640 }
641
642 return AVSValue();
643 }
644
645 void SchedulePrefetch(int currentN, InternalEnvironment* env)
646 {
647 int numQueued = 0;
648 for (int n = currentN + 1;
649 (n < currentN + prefetchFrames) && (n < vi.num_frames);
650 ++n)
651 {
652 PVideoFrame result;
653 CacheType::handle cacheHandle;
654 switch (videoCache->lookup(n, &cacheHandle, false, result))
655 {
656 case LRU_LOOKUP_NOT_FOUND:
657 {
658 std::lock_guard<std::mutex> lock(mutex);
659 prefetchQueue.emplace_back(n, cacheHandle);
660 ++numQueued;
661 break;
662 }
663 case LRU_LOOKUP_FOUND_AND_READY: // Fall-through intentional
664 case LRU_LOOKUP_NO_CACHE: // Fall-through intentional
665 case LRU_LOOKUP_FOUND_BUT_NOTAVAIL:
666 {
667 break;
668 }
669 default:
670 {
671 env->ThrowError("Invalid Program");
672 break;
673 }
674 }
675 }
676 // wake up sleeping thread
677 if (numQueued > 0) {
678 std::lock_guard<std::mutex> lock(mutex);
679 for (; numWorkers < numThreads && numQueued > 0; ++numWorkers, --numQueued) {
680 threadPool->QueueJob(ThreadWorker_, this, env, nullptr);
681 }
682 }
683 }
684
685 public:
686 ✗ QueuePrefetcher(PClip child, int prefetchFrames, int numThreads, Device* device, InternalEnvironment* env) :
687 ✗ child(child),
688 ✗ vi(child->GetVideoInfo()),
689 ✗ prefetchFrames(prefetchFrames),
690 ✗ numThreads(numThreads),
691 ✗ threadPool(NULL),
692 ✗ device(device),
693 ✗ videoCache(NULL),
694 ✗ numWorkers(0),
695 ✗ workerExceptionPresent(false)
696 {
697 ✗ if (numThreads > 0 && prefetchFrames > 0) {
698 ✗ threadPool = env->NewThreadPool(numThreads);
699 ✗ videoCache = std::shared_ptr<CacheType>(new CacheType(prefetchFrames - 1, CACHE_NO_RESIZE));
700 }
701 else {
702 ✗ numThreads = prefetchFrames = 0;
703 }
704 ✗ }
705
706 ✗ ~QueuePrefetcher()
707 {
708 ✗ if (numThreads > 0) {
709 // finish threadpool
710 ✗ threadPool->Finish();
711
712 // cancel queue
713 ✗ while (prefetchQueue.size() > 0) {
714 ✗ videoCache->rollback(&prefetchQueue.front().second);
715 ✗ prefetchQueue.pop_front();
716 }
717 }
718 ✗ }
719
720 VideoInfo GetVideoInfo() const { return vi; }
721
722 PVideoFrame GetFrame(int n, InternalEnvironment* env)
723 {
724 // do not use thread when invoke running
725 if (prefetchFrames == 0 || env->GetSuppressThreadCount() > 0) {
726 return child->GetFrame(n, env);
727 }
728
729 PVideoFrame result;
730 CacheType::handle cacheHandle;
731 switch (videoCache->lookup(n, &cacheHandle, true, result)) {
732 case LRU_LOOKUP_FOUND_AND_READY:
733 break;
734 case LRU_LOOKUP_NOT_FOUND:
735 {
736 try
737 {
738 {
739 std::lock_guard<std::mutex> lock(mutex);
740 // check error
741 if (workerExceptionPresent)
742 {
743 std::rethrow_exception(workerException);
744 }
745 }
746 cacheHandle.first->value = child->GetFrame(n, env);
747 result = cacheHandle.first->value;
748 videoCache->commit_value(&cacheHandle);
749 }
750 catch (...)
751 {
752 videoCache->rollback(&cacheHandle);
753 throw;
754 }
755 break;
756 }
757 case LRU_LOOKUP_NO_CACHE:
758 result = child->GetFrame(n, env);
759 break;
760 case LRU_LOOKUP_FOUND_BUT_NOTAVAIL: // Fall-through intentional
761 default:
762 env->ThrowError("Invalid Program");
763 break;
764 }
765
766 SchedulePrefetch(n, env);
767
768 return result;
769 }
770 };
771
772 class FrameTransferEngine
773 {
774 public:
775 QueuePrefetcher& child;
776 VideoInfo vi;
777
778 Device* upstreamDevice;
779 Device* downstreamDevice;
780
781 FrameTransferEngine(QueuePrefetcher& child, Device* upstreamDevice, Device* downstreamDevice) :
782 child(child),
783 vi(child.GetVideoInfo()),
784 upstreamDevice(upstreamDevice),
785 downstreamDevice(downstreamDevice)
786 { }
787
788 virtual ~FrameTransferEngine() { }
789
790 virtual PVideoFrame GetFrame(int n, InternalEnvironment* env) = 0;
791 };
792
793 #ifdef ENABLE_CUDA
794 class CUDAFrameTransferEngine : public FrameTransferEngine
795 {
796 typedef LruCache<size_t, PVideoFrame> CacheType;
797
798 struct QueueItem {
799 size_t n;
800 PVideoFrame src;
801 CacheType::handle cacheHandle;
802 cudaEvent_t completeEvent;
803 std::unique_ptr<std::vector<DeviceCompleteCallbackData>> completeCallbacks;
804 };
805
806 int prefetchFrames;
807
808 std::shared_ptr<CacheType> videoCache;
809
810 std::mutex mutex;
811 cudaStream_t stream;
812 std::deque<QueueItem> prefetchQueue;
813
814 cudaMemcpyKind GetMemcpyKind()
815 {
816 if (upstreamDevice->device_type == DEV_TYPE_CPU && downstreamDevice->device_type == DEV_TYPE_CUDA) {
817 // Host to Device
818 return cudaMemcpyHostToDevice;
819 }
820 if (upstreamDevice->device_type == DEV_TYPE_CUDA && downstreamDevice->device_type == DEV_TYPE_CPU) {
821 // Device to Host
822 return cudaMemcpyDeviceToHost;
823 }
824 if (upstreamDevice->device_type == DEV_TYPE_CUDA && downstreamDevice->device_type == DEV_TYPE_CUDA) {
825 // Device to Device
826 return cudaMemcpyDeviceToDevice;
827 }
828 _ASSERT(false);
829 return cudaMemcpyDefault;
830 }
831
832 void ExecuteCallbacks(const std::vector<DeviceCompleteCallbackData>* callbacks)
833 {
834 if (callbacks != nullptr) {
835 for (auto cbdata : *callbacks) {
836 cbdata.cb(cbdata.user_data);
837 }
838 }
839 }
840
841 void TransferFrameData(PVideoFrame& dst, PVideoFrame& src, bool async, InternalEnvironment* env)
842 {
843 size_t srchead = GetFrameHead(src);
844 size_t sz = GetFrameTail(src) - srchead;
845 const BYTE* srcptr = src->GetFrameBuffer()->GetReadPtr() + srchead;
846 BYTE* dstptr = dst->GetFrameBuffer()->GetWritePtr() + GetFrameHead(dst);
847 cudaMemcpyKind kind = GetMemcpyKind();
848
849 if (async) {
850 CUDA_CHECK(cudaMemcpyAsync(dstptr, srcptr, sz, kind, stream));
851 }
852 else {
853 CUDA_CHECK(cudaMemcpy(dstptr, srcptr, sz, kind));
854 }
855 }
856
857 PVideoFrame GetFrameImmediate(int n, InternalEnvironment* env)
858 {
859 PVideoFrame src = child.GetFrame(n, env);
860 PVideoFrame dst = env->GetOnDeviceFrame(src, downstreamDevice);
861 TransferFrameData(dst, src, false, env);
862
863 #if 1
864 AVSMap* mapv = env->getFramePropsRW(dst);
865 const int numKeys = env->propNumKeys(mapv);
866 for (int i = 0; i < numKeys; i++) {
867 const char* key = env->propGetKey(mapv, i);
868 if (env->propGetType(mapv, key) == AVSPropTypes::PROPTYPE_FRAME) {
869 // isFrame true
870 const int numElements = env->propNumElements(mapv, key);
871
872 std::vector<PVideoFrame> frameset;
873 int error;
874 // avs+: can be more frames in a frame property array
875 for (int index = 0; index < numElements; index++) {
876 const PVideoFrame srcframe = env->propGetFrame(mapv, key, index, &error);
877 frameset.push_back(srcframe);
878 }
879
880 env->propDeleteKey(mapv, key);
881
882 for (int index = 0; index < numElements; index++) {
883 PVideoFrame src = frameset[index];
884 PVideoFrame dst = env->GetOnDeviceFrame(src, downstreamDevice);
885 TransferFrameData(dst, src, false, env);
886 env->propSetFrame(mapv, key, dst, AVSPropAppendMode::PROPAPPENDMODE_APPEND);
887 }
888 }
889 }
890 #else
891 // kept for reference from neo fork, a single frame in frame properties. No array of frames here
892 AVSMap* mapv = env->GetAVSMap(dst);
893 for (auto it = mapv->data.begin(), end = mapv->data.end(); it != end; ++it) {
894 if (it->second.IsFrame()) {
895 PVideoFrame src = it->second.GetFrame();
896 PVideoFrame dst = env->GetOnDeviceFrame(src, downstreamDevice);
897 TransferFrameData(dst, src, false, env);
898 it->second = dst;
899 }
900 }
901 #endif
902 ExecuteCallbacks(downstreamDevice->GetAndClearCallbacks().get());
903
904 return dst;
905 }
906
907 QueueItem SetupTransfer(int n, CacheType::handle& cacheHandle, InternalEnvironment* env)
908 {
909 QueueItem item = { (size_t)n, child.GetFrame(n, env), cacheHandle, nullptr, nullptr };
910 cacheHandle.first->value = env->GetOnDeviceFrame(item.src, downstreamDevice);
911 CUDA_CHECK(cudaEventCreate(&item.completeEvent));
912
913 item.completeCallbacks = upstreamDevice->GetAndClearCallbacks();
914
915 if (upstreamDevice->device_type == DEV_TYPE_CUDA) {
916 static_cast<CUDADevice*>(upstreamDevice)->MakeStreamWaitCompute(stream, env);
917 }
918
919 TransferFrameData(cacheHandle.first->value, item.src, true, env);
920
921 #if 1
922 AVSMap* mapv = env->getFramePropsRW(cacheHandle.first->value);
923 const int numKeys = env->propNumKeys(mapv);
924 for (int i = 0; i < numKeys; i++) {
925 const char* key = env->propGetKey(mapv, i);
926 if (env->propGetType(mapv, key) == AVSPropTypes::PROPTYPE_FRAME) {
927 // isFrame true
928 const int numElements = env->propNumElements(mapv, key);
929
930 std::vector<PVideoFrame> frameset;
931 int error;
932 // avs+: can be more frames in a frame property array
933 for (int index = 0; index < numElements; index++) {
934 const PVideoFrame srcframe = env->propGetFrame(mapv, key, index, &error);
935 frameset.push_back(srcframe);
936 }
937
938 env->propDeleteKey(mapv, key);
939
940 for (int index = 0; index < numElements; index++) {
941 PVideoFrame src = frameset[index];
942 PVideoFrame dst = env->GetOnDeviceFrame(src, downstreamDevice);
943 TransferFrameData(dst, src, true, env);
944 env->propSetFrame(mapv, key, dst, AVSPropAppendMode::PROPAPPENDMODE_APPEND);
945 }
946 }
947 }
948 #else
949 // kept for reference from neo fork, a single frame in frame properties. No frame arrays here
950 AVSMap* mapv = env->GetAVSMap(cacheHandle.first->value);
951 for (auto it = mapv->data.begin(), end = mapv->data.end(); it != end; ++it) {
952 if (it->second.IsFrame()) {
953 PVideoFrame src = it->second.GetFrame();
954 PVideoFrame dst = env->GetOnDeviceFrame(src, downstreamDevice);
955 TransferFrameData(dst, src, true, env);
956 it->second = dst;
957 }
958 }
959 #endif
960
961 CUDA_CHECK(cudaEventRecord(item.completeEvent, stream));
962
963 return std::move(item);
964 }
965
966 int SchedulePrefetch(int currentN, int prefetchStart, InternalEnvironment* env)
967 {
968 int numQueued = 0;
969 int n = prefetchStart;
970 for (; n < currentN + prefetchFrames && n < vi.num_frames; ++n)
971 {
972 PVideoFrame result;
973 CacheType::handle cacheHandle;
974 switch (videoCache->lookup(n, &cacheHandle, false, result))
975 {
976 case LRU_LOOKUP_NOT_FOUND:
977 {
978 try {
979 prefetchQueue.emplace_back(SetupTransfer(n, cacheHandle, env));
980 }
981 catch (...) {
982 videoCache->rollback(&cacheHandle);
983 throw;
984 }
985 break;
986 }
987 case LRU_LOOKUP_FOUND_AND_READY: // Fall-through intentional
988 case LRU_LOOKUP_NO_CACHE: // Fall-through intentional
989 case LRU_LOOKUP_FOUND_BUT_NOTAVAIL:
990 {
991 break;
992 }
993 default:
994 {
995 env->ThrowError("Invalid Program");
996 break;
997 }
998 }
999 }
1000 return n;
1001 }
1002
1003 void FinishCompleted(InternalEnvironment* env)
1004 {
1005 while (prefetchQueue.size() > 0) {
1006 QueueItem& item = prefetchQueue.front();
1007 cudaError_t err = cudaEventQuery(item.completeEvent);
1008 if (err == cudaErrorNotReady) {
1009 break;
1010 }
1011 try {
1012 CUDA_CHECK(err);
1013
1014 // transfer is complete
1015 CUDA_CHECK(cudaEventDestroy(item.completeEvent));
1016 ExecuteCallbacks(item.completeCallbacks.get());
1017
1018 videoCache->commit_value(&item.cacheHandle);
1019 }
1020 catch (...) {
1021 videoCache->rollback(&item.cacheHandle);
1022 throw;
1023 }
1024 prefetchQueue.pop_front();
1025 }
1026 }
1027
1028 PVideoFrame WaitUntil(int n, InternalEnvironment* env)
1029 {
1030 while (prefetchQueue.size() > 0) {
1031 QueueItem& item = prefetchQueue.front();
1032 try {
1033 CUDA_CHECK(cudaEventSynchronize(item.completeEvent));
1034 CUDA_CHECK(cudaEventDestroy(item.completeEvent));
1035 ExecuteCallbacks(item.completeCallbacks.get());
1036
1037 PVideoFrame frame = item.cacheHandle.first->value; // fill before Commit !!!
1038
1039 videoCache->commit_value(&item.cacheHandle);
1040
1041 if (item.n == n) {
1042 prefetchQueue.pop_front();
1043 return frame;
1044 }
1045
1046 prefetchQueue.pop_front();
1047 }
1048 catch (...) {
1049 videoCache->rollback(&item.cacheHandle);
1050 throw;
1051 }
1052 }
1053 env->ThrowError("invalid program");
1054 return PVideoFrame();
1055 }
1056
1057 void CheckDevicePair(InternalEnvironment* env)
1058 {
1059 if (upstreamDevice->device_type == DEV_TYPE_CPU && downstreamDevice->device_type == DEV_TYPE_CUDA) {
1060 // Host to Device
1061 return;
1062 }
1063 if (upstreamDevice->device_type == DEV_TYPE_CUDA && downstreamDevice->device_type == DEV_TYPE_CPU) {
1064 // Device to Host
1065 return;
1066 }
1067 if (upstreamDevice->device_type == DEV_TYPE_CUDA && downstreamDevice->device_type == DEV_TYPE_CUDA) {
1068 // Device to Device
1069 return;
1070 }
1071 env->ThrowError("[CUDAFrameTransferEngine] invalid device pair up:%s down:%d",
1072 upstreamDevice->GetName(), downstreamDevice->GetName());
1073 }
1074
1075 public:
1076 CUDAFrameTransferEngine(QueuePrefetcher& child, Device* upstreamDevice, Device* downstreamDevice, int prefetchFrames, InternalEnvironment* env) :
1077 FrameTransferEngine(child, upstreamDevice, downstreamDevice),
1078 prefetchFrames(prefetchFrames),
1079 videoCache(new CacheType(std::max(0, prefetchFrames - 1), CACHE_NO_RESIZE)),
1080 stream(nullptr)
1081 {
1082 CheckDevicePair(env);
1083
1084 // note: stream belongs to a device
1085 upstreamDevice->SetActiveToCurrentThread(env);
1086 CUDA_CHECK(cudaStreamCreate(&stream));
1087 }
1088
1089 ~CUDAFrameTransferEngine()
1090 {
1091 for (auto& item : prefetchQueue) {
1092 cudaEventSynchronize(item.completeEvent);
1093 cudaEventDestroy(item.completeEvent);
1094 ExecuteCallbacks(item.completeCallbacks.get());
1095 videoCache->commit_value(&item.cacheHandle);
1096 }
1097 prefetchQueue.clear();
1098
1099 cudaStreamDestroy(stream);
1100 }
1101
1102 virtual PVideoFrame GetFrame(int n, InternalEnvironment* env)
1103 {
1104 // Giant lock. This is OK because all transfer is done asynchronously
1105 std::lock_guard<std::mutex> lock(mutex);
1106
1107 // set upstream device
1108 upstreamDevice->SetActiveToCurrentThread(env);
1109
1110 // do not use thread when invoke running
1111 if (prefetchFrames == 0 || env->GetSuppressThreadCount() > 0) {
1112 return GetFrameImmediate(n, env);
1113 }
1114
1115 FinishCompleted(env);
1116
1117 SchedulePrefetch(n, n, env);
1118
1119 // Get requested frame
1120 PVideoFrame result;
1121 CacheType::handle cacheHandle;
1122 // fill result if LRU_LOOKUP_FOUND_AND_READY
1123 switch (videoCache->lookup(n, &cacheHandle, false, result))
1124 {
1125 case LRU_LOOKUP_FOUND_AND_READY:
1126 {
1127 break;
1128 }
1129 case LRU_LOOKUP_NO_CACHE:
1130 {
1131 result = GetFrameImmediate(n, env);
1132 break;
1133 }
1134 case LRU_LOOKUP_FOUND_BUT_NOTAVAIL:
1135 {
1136 // now transferring, wait until completion
1137 result = WaitUntil(n, env);
1138 break;
1139 }
1140 case LRU_LOOKUP_NOT_FOUND:
1141 {
1142 env->ThrowError("Oh.. maybe cache size is too small ...");
1143 break;
1144 }
1145 default:
1146 {
1147 env->ThrowError("Invalid Program");
1148 break;
1149 }
1150 }
1151
1152 // set downstreamdevice
1153 downstreamDevice->SetActiveToCurrentThread(env);
1154
1155 return result;
1156 }
1157 };
1158 #endif
1159
1160 ✗ FrameTransferEngine* CreateTransferEngine(QueuePrefetcher& child,
1161 Device* upstreamDevice, Device* downstreamDevice, int prefetchFrames, InternalEnvironment* env)
1162 {
1163 #ifdef ENABLE_CUDA
1164 if (upstreamDevice->device_type == DEV_TYPE_CPU && downstreamDevice->device_type == DEV_TYPE_CUDA) {
1165 // CPU to CUDA
1166 return new CUDAFrameTransferEngine(child, upstreamDevice, downstreamDevice, prefetchFrames, env);
1167 }
1168 if (upstreamDevice->device_type == DEV_TYPE_CUDA && downstreamDevice->device_type == DEV_TYPE_CPU) {
1169 // CUDA to CPU
1170 return new CUDAFrameTransferEngine(child, upstreamDevice, downstreamDevice, prefetchFrames, env);
1171 }
1172 if (upstreamDevice->device_type == DEV_TYPE_CUDA && downstreamDevice->device_type == DEV_TYPE_CUDA) {
1173 // CUDA to CUDA
1174 return new CUDAFrameTransferEngine(child, upstreamDevice, downstreamDevice, prefetchFrames, env);
1175 }
1176 #endif
1177 ✗ env->ThrowError("Not supported frame data transfer. up:%s down:%d",
1178 ✗ upstreamDevice->GetName(), downstreamDevice->GetName());
1179 ✗ return nullptr;
1180 }
1181
1182 class OnDevice : public GenericVideoFilter
1183 {
1184 Device* upstreamDevice;
1185 int prefetchFrames;
1186
1187 QueuePrefetcher prefetcher;
1188
1189 std::mutex mutex;
1190 std::map<Device*, std::unique_ptr<FrameTransferEngine>> transferEngines;
1191
1192 ✗ FrameTransferEngine* GetOrCreateTransferEngine(Device* downstreamDevice, InternalEnvironment* env)
1193 {
1194 ✗ std::lock_guard<std::mutex> lock(mutex);
1195 ✗ auto it = transferEngines.find(downstreamDevice);
1196 ✗ if (it != transferEngines.end()) {
1197 ✗ return it->second.get();
1198 }
1199 ✗ int transferPrefetch = (prefetchFrames == 1) ? 2 : prefetchFrames;
1200 ✗ auto pEngine = CreateTransferEngine(prefetcher, upstreamDevice, downstreamDevice, transferPrefetch, env);
1201 ✗ transferEngines[downstreamDevice] = std::unique_ptr<FrameTransferEngine>(pEngine);
1202 ✗ return pEngine;
1203 ✗ }
1204
1205 public:
1206 ✗ OnDevice(PClip child, int prefetchFrames, Device* upstreamDevice, InternalEnvironment* env) :
1207 GenericVideoFilter(child),
1208 ✗ upstreamDevice(upstreamDevice),
1209 ✗ prefetchFrames(prefetchFrames),
1210 ✗ prefetcher(child, (prefetchFrames >= 2) ? 2 : 0, (prefetchFrames >= 2) ? 1 : 0, upstreamDevice, env)
1211 ✗ { }
1212
1213 ✗ PVideoFrame __stdcall GetFrame(int n, IScriptEnvironment* env_)
1214 {
1215 ✗ InternalEnvironment* env = GetAndRevealCamouflagedEnv(env_);
1216 ✗ Device* downstreamDevice = env->SetCurrentDevice(upstreamDevice);
1217
1218 ✗ if (downstreamDevice == nullptr) {
1219 ✗ env->ThrowError("This thread is not created by AviSynth. It is not allowed to call GetFrame on this thread ...");
1220 }
1221
1222 ✗ if (downstreamDevice == upstreamDevice) {
1223 // shortcut
1224 ✗ return child->GetFrame(n, env);
1225 }
1226
1227 // Get frame via transfer engine
1228 ✗ PVideoFrame frame = GetOrCreateTransferEngine(downstreamDevice, env)->GetFrame(n, env);
1229
1230 ✗ env->SetCurrentDevice(downstreamDevice);
1231 ✗ return frame;
1232 ✗ }
1233
1234 ✗ void __stdcall GetAudio(void* buf, int64_t start, int64_t count, IScriptEnvironment* env_)
1235 {
1236 ✗ InternalEnvironment* env = GetAndRevealCamouflagedEnv(env_);
1237 ✗ Device* downstreamDevice = env->SetCurrentDevice(upstreamDevice);
1238 try {
1239 ✗ child->GetAudio(buf, start, count, env);
1240 ✗ env->SetCurrentDevice(downstreamDevice);
1241 }
1242 ✗ catch (...) {
1243 ✗ env->SetCurrentDevice(downstreamDevice);
1244 ✗ throw;
1245 ✗ }
1246 ✗ }
1247
1248 ✗ int __stdcall SetCacheHints(int cachehints, int frame_range)
1249 {
1250 ✗ if (cachehints == CACHE_GET_MTMODE) {
1251 ✗ return MT_NICE_FILTER;
1252 }
1253 ✗ if (cachehints == CACHE_GET_DEV_TYPE) {
1254 ✗ return DEV_TYPE_ANY;
1255 }
1256 ✗ if (cachehints == CACHE_GET_CHILD_DEV_TYPE) {
1257 ✗ return upstreamDevice->device_type;
1258 }
1259 ✗ return 0;
1260 }
1261
1262 ✗ static AVSValue __cdecl Create(AVSValue args, void* user_data, IScriptEnvironment* env_)
1263 {
1264 ✗ AvsDeviceType upstreamType = (AvsDeviceType)(size_t)user_data;
1265 ✗ InternalEnvironment* env = static_cast<InternalEnvironment*>(env_);
1266
1267 ✗ if (args[0].IsClip()) {
1268 ✗ PClip clip = args[0].AsClip();
1269 ✗ int numPrefetch = args[1].Defined() ? args[1].AsInt() : 1;
1270 ✗ int upstreamIndex = (args.ArraySize() >= 3 && args[2].Defined()) ? args[2].AsInt() : 0;
1271
1272 ✗ if (numPrefetch < 0) {
1273 ✗ numPrefetch = 0;
1274 }
1275
1276 ✗ switch (upstreamType) {
1277 ✗ case DEV_TYPE_CPU:
1278 ✗ return new OnDevice(clip, numPrefetch, (Device*)(void*)env->GetDevice(DEV_TYPE_CPU, 0), env);
1279 ✗ case DEV_TYPE_CUDA:
1280 ✗ return new OnDevice(clip, numPrefetch, (Device*)(void*)env->GetDevice(DEV_TYPE_CUDA, upstreamIndex), env);
1281 ✗ default:
1282 ✗ env->ThrowError("Not supported device ...");
1283 ✗ return AVSValue();
1284 }
1285
1286 ✗ }
1287 else {
1288 ✗ assert(args[0].IsFunction());
1289 ✗ PFunction func = args[0].AsFunction();
1290 ✗ int upstreamIndex = (args.ArraySize() >= 2 && args[1].Defined()) ? args[1].AsInt() : 0;
1291
1292 ✗ Device* upstreamDevice = nullptr;
1293 ✗ switch (upstreamType) {
1294 ✗ case DEV_TYPE_CPU:
1295 ✗ upstreamDevice = (Device*)(void*)env->GetDevice(DEV_TYPE_CPU, 0);
1296 ✗ break;
1297 ✗ case DEV_TYPE_CUDA:
1298 ✗ upstreamDevice = (Device*)(void*)env->GetDevice(DEV_TYPE_CUDA, upstreamIndex);
1299 ✗ break;
1300 ✗ default:
1301 ✗ env->ThrowError("Not supported device ...");
1302 ✗ break;
1303 }
1304
1305 ✗ DeviceSetter setter(env, upstreamDevice);
1306
1307 try {
1308 ✗ AVSValue ret = env->Invoke3(AVSValue(), func, AVSValue(nullptr, 0));
1309 ✗ return ret;
1310 ✗ }
1311 ✗ catch (IScriptEnvironment::NotFound) {
1312 ✗ const char* name = (upstreamType == DEV_TYPE_CPU) ? "OnCPU" : "OnCUDA";
1313 ✗ env->ThrowError(
1314 "%s: Invalid function parameter type '%s'(%s)\n"
1315 "Function should have no argument",
1316 ✗ name, func->GetDefinition()->param_types, func->ToString(env));
1317 ✗ }
1318
1319 ✗ return AVSValue();
1320 ✗ }
1321 }
1322 };
1323
1324 ✗ void CopyCUDAFrame(const PVideoFrame& dst, const PVideoFrame& src, InternalEnvironment* env, bool sync)
1325 {
1326 #ifdef ENABLE_CUDA
1327 size_t srchead = GetFrameHead(src);
1328 size_t sz = GetFrameTail(src) - srchead;
1329 const BYTE* srcptr = src->GetFrameBuffer()->GetReadPtr() + srchead;
1330 BYTE* dstptr = dst->GetFrameBuffer()->GetWritePtr() + GetFrameHead(dst);
1331
1332 AvsDeviceType srcDevice = src->GetDevice()->device_type;
1333 AvsDeviceType dstDevice = dst->GetDevice()->device_type;
1334 cudaMemcpyKind kind = cudaMemcpyHostToHost;
1335
1336 if (srcDevice == DEV_TYPE_CPU && dstDevice == DEV_TYPE_CUDA) {
1337 kind = cudaMemcpyHostToDevice;
1338 }
1339 if (srcDevice == DEV_TYPE_CUDA && dstDevice == DEV_TYPE_CPU) {
1340 kind = cudaMemcpyDeviceToHost;
1341 }
1342 if (srcDevice == DEV_TYPE_CUDA && dstDevice == DEV_TYPE_CUDA) {
1343 kind = cudaMemcpyDeviceToDevice;
1344 }
1345
1346 if (sync) {
1347 CUDA_CHECK(cudaMemcpy(dstptr, srcptr, sz, kind));
1348 }
1349 else {
1350 CUDA_CHECK(cudaMemcpyAsync(dstptr, srcptr, sz, kind));
1351 }
1352 #else
1353 ✗ env->ThrowError("CopyCUDAFrame: CUDA support is disabled ...");
1354 #endif
1355 ✗ }
1356
1357 ✗ PVideoFrame GetFrameOnDevice(PClip& c, int n, const PDevice& device, InternalEnvironment* env)
1358 {
1359 ✗ DeviceSetter setter(env, (Device*)(void*)device);
1360 ✗ return c->GetFrame(n, env);
1361 ✗ }
1362
1363 extern const AVSFunction Device_filters[] = {
1364 { "OnCPU", BUILTIN_FUNC_PREFIX, "c[num_prefetch]i", OnDevice::Create, (void*)DEV_TYPE_CPU },
1365 { "OnCUDA", BUILTIN_FUNC_PREFIX, "c[num_prefetch]i[device_index]i", OnDevice::Create, (void*)DEV_TYPE_CUDA },
1366 { "OnCPU", BUILTIN_FUNC_PREFIX, "n", OnDevice::Create, (void*)DEV_TYPE_CPU },
1367 { "OnCUDA", BUILTIN_FUNC_PREFIX, "n[device_index]i", OnDevice::Create, (void*)DEV_TYPE_CUDA },
1368 { 0 }
1369 };
1370