← Back to branch summary

~paparazzi-uav/paparazzi/v5.0-manual

« back to all changes in this revision

Viewing changes to sw/ext/opencv_bebop/opencv/modules/core/src/cuda_stream.cpp

  • Committer: Paparazzi buildbot
  • Date: 2016-05-18 15:00:29 UTC
  • Revision ID: felix.ruess+docbot@gmail.com-20160518150029-e8lgzi5kvb4p7un9
Manual import commit 4b8bbb730080dac23cf816b98908dacfabe2a8ec from v5.0 branch.

Show diffs side-by-side

added added

removed removed

Lines of Context:
sw/ext/opencv_bebop/opencv/modules/core/src/cuda_stream.cpp
 
1
/*M///////////////////////////////////////////////////////////////////////////////////////
 
2
//
 
3
//  IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
 
4
//
 
5
//  By downloading, copying, installing or using the software you agree to this license.
 
6
//  If you do not agree to this license, do not download, install,
 
7
//  copy or use the software.
 
8
//
 
9
//
 
10
//                           License Agreement
 
11
//                For Open Source Computer Vision Library
 
12
//
 
13
// Copyright (C) 2000-2008, Intel Corporation, all rights reserved.
 
14
// Copyright (C) 2009, Willow Garage Inc., all rights reserved.
 
15
// Third party copyrights are property of their respective owners.
 
16
//
 
17
// Redistribution and use in source and binary forms, with or without modification,
 
18
// are permitted provided that the following conditions are met:
 
19
//
 
20
//   * Redistribution's of source code must retain the above copyright notice,
 
21
//     this list of conditions and the following disclaimer.
 
22
//
 
23
//   * Redistribution's in binary form must reproduce the above copyright notice,
 
24
//     this list of conditions and the following disclaimer in the documentation
 
25
//     and/or other materials provided with the distribution.
 
26
//
 
27
//   * The name of the copyright holders may not be used to endorse or promote products
 
28
//     derived from this software without specific prior written permission.
 
29
//
 
30
// This software is provided by the copyright holders and contributors "as is" and
 
31
// any express or implied warranties, including, but not limited to, the implied
 
32
// warranties of merchantability and fitness for a particular purpose are disclaimed.
 
33
// In no event shall the Intel Corporation or contributors be liable for any direct,
 
34
// indirect, incidental, special, exemplary, or consequential damages
 
35
// (including, but not limited to, procurement of substitute goods or services;
 
36
// loss of use, data, or profits; or business interruption) however caused
 
37
// and on any theory of liability, whether in contract, strict liability,
 
38
// or tort (including negligence or otherwise) arising in any way out of
 
39
// the use of this software, even if advised of the possibility of such damage.
 
40
//
 
41
//M*/
 
42
 
 
43
#include "precomp.hpp"
 
44
 
 
45
using namespace cv;
 
46
using namespace cv::cuda;
 
47
 
 
48
/////////////////////////////////////////////////////////////
 
49
/// MemoryStack
 
50
 
 
51
#ifdef HAVE_CUDA
 
52
 
 
53
namespace
 
54
{
 
55
    class MemoryPool;
 
56
 
 
57
    class MemoryStack
 
58
    {
 
59
    public:
 
60
        uchar* requestMemory(size_t size);
 
61
        void returnMemory(uchar* ptr);
 
62
 
 
63
        uchar* datastart;
 
64
        uchar* dataend;
 
65
        uchar* tip;
 
66
 
 
67
        bool isFree;
 
68
        MemoryPool* pool;
 
69
 
 
70
    #if !defined(NDEBUG)
 
71
        std::vector<size_t> allocations;
 
72
    #endif
 
73
    };
 
74
 
 
75
    uchar* MemoryStack::requestMemory(size_t size)
 
76
    {
 
77
        const size_t freeMem = dataend - tip;
 
78
 
 
79
        if (size > freeMem)
 
80
            return 0;
 
81
 
 
82
        uchar* ptr = tip;
 
83
 
 
84
        tip += size;
 
85
 
 
86
    #if !defined(NDEBUG)
 
87
        allocations.push_back(size);
 
88
    #endif
 
89
 
 
90
        return ptr;
 
91
    }
 
92
 
 
93
    void MemoryStack::returnMemory(uchar* ptr)
 
94
    {
 
95
        CV_DbgAssert( ptr >= datastart && ptr < dataend );
 
96
 
 
97
    #if !defined(NDEBUG)
 
98
        const size_t allocSize = tip - ptr;
 
99
        CV_Assert( allocSize == allocations.back() );
 
100
        allocations.pop_back();
 
101
    #endif
 
102
 
 
103
        tip = ptr;
 
104
    }
 
105
}
 
106
 
 
107
#endif
 
108
 
 
109
/////////////////////////////////////////////////////////////
 
110
/// MemoryPool
 
111
 
 
112
#ifdef HAVE_CUDA
 
113
 
 
114
namespace
 
115
{
 
116
    class MemoryPool
 
117
    {
 
118
    public:
 
119
        MemoryPool();
 
120
 
 
121
        void initialize(size_t stackSize, int stackCount);
 
122
        void release();
 
123
 
 
124
        MemoryStack* getFreeMemStack();
 
125
        void returnMemStack(MemoryStack* memStack);
 
126
 
 
127
    private:
 
128
        void initilizeImpl();
 
129
 
 
130
        Mutex mtx_;
 
131
 
 
132
        bool initialized_;
 
133
        size_t stackSize_;
 
134
        int stackCount_;
 
135
 
 
136
        uchar* mem_;
 
137
 
 
138
        std::vector<MemoryStack> stacks_;
 
139
    };
 
140
 
 
141
    MemoryPool::MemoryPool() : initialized_(false), mem_(0)
 
142
    {
 
143
        // default : 10 Mb, 5 stacks
 
144
        stackSize_ = 10 * 1024 * 1024;
 
145
        stackCount_ = 5;
 
146
    }
 
147
 
 
148
    void MemoryPool::initialize(size_t stackSize, int stackCount)
 
149
    {
 
150
        AutoLock lock(mtx_);
 
151
 
 
152
        release();
 
153
 
 
154
        stackSize_ = stackSize;
 
155
        stackCount_ = stackCount;
 
156
 
 
157
        initilizeImpl();
 
158
    }
 
159
 
 
160
    void MemoryPool::initilizeImpl()
 
161
    {
 
162
        const size_t totalSize = stackSize_ * stackCount_;
 
163
 
 
164
        if (totalSize > 0)
 
165
        {
 
166
            cudaError_t err = cudaMalloc(&mem_, totalSize);
 
167
            if (err != cudaSuccess)
 
168
                return;
 
169
 
 
170
            stacks_.resize(stackCount_);
 
171
 
 
172
            uchar* ptr = mem_;
 
173
 
 
174
            for (int i = 0; i < stackCount_; ++i)
 
175
            {
 
176
                stacks_[i].datastart = ptr;
 
177
                stacks_[i].dataend = ptr + stackSize_;
 
178
                stacks_[i].tip = ptr;
 
179
                stacks_[i].isFree = true;
 
180
                stacks_[i].pool = this;
 
181
 
 
182
                ptr += stackSize_;
 
183
            }
 
184
 
 
185
            initialized_ = true;
 
186
        }
 
187
    }
 
188
 
 
189
    void MemoryPool::release()
 
190
    {
 
191
        if (mem_)
 
192
        {
 
193
#if !defined(NDEBUG)
 
194
            for (int i = 0; i < stackCount_; ++i)
 
195
            {
 
196
                CV_DbgAssert( stacks_[i].isFree );
 
197
                CV_DbgAssert( stacks_[i].tip == stacks_[i].datastart );
 
198
            }
 
199
#endif
 
200
 
 
201
            cudaFree(mem_);
 
202
 
 
203
            mem_ = 0;
 
204
            initialized_ = false;
 
205
        }
 
206
    }
 
207
 
 
208
    MemoryStack* MemoryPool::getFreeMemStack()
 
209
    {
 
210
        AutoLock lock(mtx_);
 
211
 
 
212
        if (!initialized_)
 
213
            initilizeImpl();
 
214
 
 
215
        if (!mem_)
 
216
            return 0;
 
217
 
 
218
        for (int i = 0; i < stackCount_; ++i)
 
219
        {
 
220
            if (stacks_[i].isFree)
 
221
            {
 
222
                stacks_[i].isFree = false;
 
223
                return &stacks_[i];
 
224
            }
 
225
        }
 
226
 
 
227
        return 0;
 
228
    }
 
229
 
 
230
    void MemoryPool::returnMemStack(MemoryStack* memStack)
 
231
    {
 
232
        AutoLock lock(mtx_);
 
233
 
 
234
        CV_DbgAssert( !memStack->isFree );
 
235
 
 
236
#if !defined(NDEBUG)
 
237
        bool found = false;
 
238
        for (int i = 0; i < stackCount_; ++i)
 
239
        {
 
240
            if (memStack == &stacks_[i])
 
241
            {
 
242
                found = true;
 
243
                break;
 
244
            }
 
245
        }
 
246
        CV_DbgAssert( found );
 
247
#endif
 
248
 
 
249
        CV_DbgAssert( memStack->tip == memStack->datastart );
 
250
 
 
251
        memStack->isFree = true;
 
252
    }
 
253
}
 
254
 
 
255
#endif
 
256
 
 
257
////////////////////////////////////////////////////////////////
 
258
/// Stream::Impl
 
259
 
 
260
#ifndef HAVE_CUDA
 
261
 
 
262
class cv::cuda::Stream::Impl
 
263
{
 
264
public:
 
265
    Impl(void* ptr = 0)
 
266
    {
 
267
        (void) ptr;
 
268
        throw_no_cuda();
 
269
    }
 
270
};
 
271
 
 
272
#else
 
273
 
 
274
namespace
 
275
{
 
276
    class StackAllocator;
 
277
}
 
278
 
 
279
class cv::cuda::Stream::Impl
 
280
{
 
281
public:
 
282
    cudaStream_t stream;
 
283
    bool ownStream;
 
284
 
 
285
    Ptr<StackAllocator> stackAllocator;
 
286
 
 
287
    Impl();
 
288
    explicit Impl(cudaStream_t stream);
 
289
 
 
290
    ~Impl();
 
291
};
 
292
 
 
293
cv::cuda::Stream::Impl::Impl() : stream(0), ownStream(false)
 
294
{
 
295
    cudaSafeCall( cudaStreamCreate(&stream) );
 
296
    ownStream = true;
 
297
 
 
298
    stackAllocator = makePtr<StackAllocator>(stream);
 
299
}
 
300
 
 
301
cv::cuda::Stream::Impl::Impl(cudaStream_t stream_) : stream(stream_), ownStream(false)
 
302
{
 
303
    stackAllocator = makePtr<StackAllocator>(stream);
 
304
}
 
305
 
 
306
cv::cuda::Stream::Impl::~Impl()
 
307
{
 
308
    stackAllocator.release();
 
309
 
 
310
    if (stream && ownStream)
 
311
    {
 
312
        cudaStreamDestroy(stream);
 
313
    }
 
314
}
 
315
 
 
316
#endif
 
317
 
 
318
/////////////////////////////////////////////////////////////
 
319
/// DefaultDeviceInitializer
 
320
 
 
321
#ifdef HAVE_CUDA
 
322
 
 
323
namespace cv { namespace cuda
 
324
{
 
325
    class DefaultDeviceInitializer
 
326
    {
 
327
    public:
 
328
        DefaultDeviceInitializer();
 
329
        ~DefaultDeviceInitializer();
 
330
 
 
331
        Stream& getNullStream(int deviceId);
 
332
        MemoryPool* getMemoryPool(int deviceId);
 
333
 
 
334
    private:
 
335
        void initStreams();
 
336
        void initPools();
 
337
 
 
338
        std::vector<Ptr<Stream> > streams_;
 
339
        Mutex streams_mtx_;
 
340
 
 
341
        std::vector<MemoryPool> pools_;
 
342
        Mutex pools_mtx_;
 
343
    };
 
344
 
 
345
    DefaultDeviceInitializer::DefaultDeviceInitializer()
 
346
    {
 
347
    }
 
348
 
 
349
    DefaultDeviceInitializer::~DefaultDeviceInitializer()
 
350
    {
 
351
        streams_.clear();
 
352
 
 
353
        for (size_t i = 0; i < pools_.size(); ++i)
 
354
        {
 
355
            cudaSetDevice(static_cast<int>(i));
 
356
            pools_[i].release();
 
357
        }
 
358
 
 
359
        pools_.clear();
 
360
    }
 
361
 
 
362
    Stream& DefaultDeviceInitializer::getNullStream(int deviceId)
 
363
    {
 
364
        AutoLock lock(streams_mtx_);
 
365
 
 
366
        if (streams_.empty())
 
367
        {
 
368
            int deviceCount = getCudaEnabledDeviceCount();
 
369
 
 
370
            if (deviceCount > 0)
 
371
                streams_.resize(deviceCount);
 
372
        }
 
373
 
 
374
        CV_DbgAssert( deviceId >= 0 && deviceId < static_cast<int>(streams_.size()) );
 
375
 
 
376
        if (streams_[deviceId].empty())
 
377
        {
 
378
            cudaStream_t stream = NULL;
 
379
            Ptr<Stream::Impl> impl = makePtr<Stream::Impl>(stream);
 
380
            streams_[deviceId] = Ptr<Stream>(new Stream(impl));
 
381
        }
 
382
 
 
383
        return *streams_[deviceId];
 
384
    }
 
385
 
 
386
    MemoryPool* DefaultDeviceInitializer::getMemoryPool(int deviceId)
 
387
    {
 
388
        AutoLock lock(pools_mtx_);
 
389
 
 
390
        if (pools_.empty())
 
391
        {
 
392
            int deviceCount = getCudaEnabledDeviceCount();
 
393
 
 
394
            if (deviceCount > 0)
 
395
                pools_.resize(deviceCount);
 
396
        }
 
397
 
 
398
        CV_DbgAssert( deviceId >= 0 && deviceId < static_cast<int>(pools_.size()) );
 
399
 
 
400
        return &pools_[deviceId];
 
401
    }
 
402
 
 
403
    DefaultDeviceInitializer initializer;
 
404
}}
 
405
 
 
406
#endif
 
407
 
 
408
/////////////////////////////////////////////////////////////
 
409
/// Stream
 
410
 
 
411
cv::cuda::Stream::Stream()
 
412
{
 
413
#ifndef HAVE_CUDA
 
414
    throw_no_cuda();
 
415
#else
 
416
    impl_ = makePtr<Impl>();
 
417
#endif
 
418
}
 
419
 
 
420
bool cv::cuda::Stream::queryIfComplete() const
 
421
{
 
422
#ifndef HAVE_CUDA
 
423
    throw_no_cuda();
 
424
    return false;
 
425
#else
 
426
    cudaError_t err = cudaStreamQuery(impl_->stream);
 
427
 
 
428
    if (err == cudaErrorNotReady || err == cudaSuccess)
 
429
        return err == cudaSuccess;
 
430
 
 
431
    cudaSafeCall(err);
 
432
    return false;
 
433
#endif
 
434
}
 
435
 
 
436
void cv::cuda::Stream::waitForCompletion()
 
437
{
 
438
#ifndef HAVE_CUDA
 
439
    throw_no_cuda();
 
440
#else
 
441
    cudaSafeCall( cudaStreamSynchronize(impl_->stream) );
 
442
#endif
 
443
}
 
444
 
 
445
void cv::cuda::Stream::waitEvent(const Event& event)
 
446
{
 
447
#ifndef HAVE_CUDA
 
448
    (void) event;
 
449
    throw_no_cuda();
 
450
#else
 
451
    cudaSafeCall( cudaStreamWaitEvent(impl_->stream, EventAccessor::getEvent(event), 0) );
 
452
#endif
 
453
}
 
454
 
 
455
#if defined(HAVE_CUDA) && (CUDART_VERSION >= 5000)
 
456
 
 
457
namespace
 
458
{
 
459
    struct CallbackData
 
460
    {
 
461
        Stream::StreamCallback callback;
 
462
        void* userData;
 
463
 
 
464
        CallbackData(Stream::StreamCallback callback_, void* userData_) : callback(callback_), userData(userData_) {}
 
465
    };
 
466
 
 
467
    void CUDART_CB cudaStreamCallback(cudaStream_t, cudaError_t status, void* userData)
 
468
    {
 
469
        CallbackData* data = reinterpret_cast<CallbackData*>(userData);
 
470
        data->callback(static_cast<int>(status), data->userData);
 
471
        delete data;
 
472
    }
 
473
}
 
474
 
 
475
#endif
 
476
 
 
477
void cv::cuda::Stream::enqueueHostCallback(StreamCallback callback, void* userData)
 
478
{
 
479
#ifndef HAVE_CUDA
 
480
    (void) callback;
 
481
    (void) userData;
 
482
    throw_no_cuda();
 
483
#else
 
484
    #if CUDART_VERSION < 5000
 
485
        (void) callback;
 
486
        (void) userData;
 
487
        CV_Error(cv::Error::StsNotImplemented, "This function requires CUDA >= 5.0");
 
488
    #else
 
489
        CallbackData* data = new CallbackData(callback, userData);
 
490
 
 
491
        cudaSafeCall( cudaStreamAddCallback(impl_->stream, cudaStreamCallback, data, 0) );
 
492
    #endif
 
493
#endif
 
494
}
 
495
 
 
496
Stream& cv::cuda::Stream::Null()
 
497
{
 
498
#ifndef HAVE_CUDA
 
499
    throw_no_cuda();
 
500
    static Stream stream;
 
501
    return stream;
 
502
#else
 
503
    const int deviceId = getDevice();
 
504
    return initializer.getNullStream(deviceId);
 
505
#endif
 
506
}
 
507
 
 
508
cv::cuda::Stream::operator bool_type() const
 
509
{
 
510
#ifndef HAVE_CUDA
 
511
    return 0;
 
512
#else
 
513
    return (impl_->stream != 0) ? &Stream::this_type_does_not_support_comparisons : 0;
 
514
#endif
 
515
}
 
516
 
 
517
#ifdef HAVE_CUDA
 
518
 
 
519
cudaStream_t cv::cuda::StreamAccessor::getStream(const Stream& stream)
 
520
{
 
521
    return stream.impl_->stream;
 
522
}
 
523
 
 
524
Stream cv::cuda::StreamAccessor::wrapStream(cudaStream_t stream)
 
525
{
 
526
    return Stream(makePtr<Stream::Impl>(stream));
 
527
}
 
528
 
 
529
#endif
 
530
 
 
531
/////////////////////////////////////////////////////////////
 
532
/// StackAllocator
 
533
 
 
534
#ifdef HAVE_CUDA
 
535
 
 
536
namespace
 
537
{
 
538
    bool enableMemoryPool = true;
 
539
 
 
540
    class StackAllocator : public GpuMat::Allocator
 
541
    {
 
542
    public:
 
543
        explicit StackAllocator(cudaStream_t stream);
 
544
        ~StackAllocator();
 
545
 
 
546
        bool allocate(GpuMat* mat, int rows, int cols, size_t elemSize);
 
547
        void free(GpuMat* mat);
 
548
 
 
549
    private:
 
550
        StackAllocator(const StackAllocator&);
 
551
        StackAllocator& operator =(const StackAllocator&);
 
552
 
 
553
        cudaStream_t stream_;
 
554
        MemoryStack* memStack_;
 
555
        size_t alignment_;
 
556
    };
 
557
 
 
558
    StackAllocator::StackAllocator(cudaStream_t stream) : stream_(stream), memStack_(0)
 
559
    {
 
560
        if (enableMemoryPool)
 
561
        {
 
562
            const int deviceId = getDevice();
 
563
            memStack_ = initializer.getMemoryPool(deviceId)->getFreeMemStack();
 
564
            DeviceInfo devInfo(deviceId);
 
565
            alignment_ = devInfo.textureAlignment();
 
566
        }
 
567
    }
 
568
 
 
569
    StackAllocator::~StackAllocator()
 
570
    {
 
571
        cudaStreamSynchronize(stream_);
 
572
 
 
573
        if (memStack_ != 0)
 
574
            memStack_->pool->returnMemStack(memStack_);
 
575
    }
 
576
 
 
577
    size_t alignUp(size_t what, size_t alignment)
 
578
    {
 
579
        size_t alignMask = alignment-1;
 
580
        size_t inverseAlignMask = ~alignMask;
 
581
        size_t res = (what + alignMask) & inverseAlignMask;
 
582
        return res;
 
583
    }
 
584
 
 
585
    bool StackAllocator::allocate(GpuMat* mat, int rows, int cols, size_t elemSize)
 
586
    {
 
587
        if (memStack_ == 0)
 
588
            return false;
 
589
 
 
590
        size_t pitch, memSize;
 
591
 
 
592
        if (rows > 1 && cols > 1)
 
593
        {
 
594
            pitch = alignUp(cols * elemSize, alignment_);
 
595
            memSize = pitch * rows;
 
596
        }
 
597
        else
 
598
        {
 
599
            // Single row or single column must be continuous
 
600
            pitch = elemSize * cols;
 
601
            memSize = alignUp(elemSize * cols * rows, 64);
 
602
        }
 
603
 
 
604
        uchar* ptr = memStack_->requestMemory(memSize);
 
605
 
 
606
        if (ptr == 0)
 
607
            return false;
 
608
 
 
609
        mat->data = ptr;
 
610
        mat->step = pitch;
 
611
        mat->refcount = (int*) fastMalloc(sizeof(int));
 
612
 
 
613
        return true;
 
614
    }
 
615
 
 
616
    void StackAllocator::free(GpuMat* mat)
 
617
    {
 
618
        if (memStack_ == 0)
 
619
            return;
 
620
 
 
621
        memStack_->returnMemory(mat->datastart);
 
622
        fastFree(mat->refcount);
 
623
    }
 
624
}
 
625
 
 
626
#endif
 
627
 
 
628
/////////////////////////////////////////////////////////////
 
629
/// BufferPool
 
630
 
 
631
void cv::cuda::setBufferPoolUsage(bool on)
 
632
{
 
633
#ifndef HAVE_CUDA
 
634
    (void)on;
 
635
    throw_no_cuda();
 
636
#else
 
637
    enableMemoryPool = on;
 
638
#endif
 
639
}
 
640
 
 
641
void cv::cuda::setBufferPoolConfig(int deviceId, size_t stackSize, int stackCount)
 
642
{
 
643
#ifndef HAVE_CUDA
 
644
    (void)deviceId;
 
645
    (void)stackSize;
 
646
    (void)stackCount;
 
647
    throw_no_cuda();
 
648
#else
 
649
    const int currentDevice = getDevice();
 
650
 
 
651
    if (deviceId >= 0)
 
652
    {
 
653
        setDevice(deviceId);
 
654
        initializer.getMemoryPool(deviceId)->initialize(stackSize, stackCount);
 
655
    }
 
656
    else
 
657
    {
 
658
        const int deviceCount = getCudaEnabledDeviceCount();
 
659
 
 
660
        for (deviceId = 0; deviceId < deviceCount; ++deviceId)
 
661
        {
 
662
            setDevice(deviceId);
 
663
            initializer.getMemoryPool(deviceId)->initialize(stackSize, stackCount);
 
664
        }
 
665
    }
 
666
 
 
667
    setDevice(currentDevice);
 
668
#endif
 
669
}
 
670
 
 
671
#ifdef HAVE_CUDA
 
672
 
 
673
cv::cuda::BufferPool::BufferPool(Stream& stream) : allocator_(stream.impl_->stackAllocator.get())
 
674
{
 
675
}
 
676
 
 
677
GpuMat cv::cuda::BufferPool::getBuffer(int rows, int cols, int type)
 
678
{
 
679
    GpuMat buf(allocator_);
 
680
    buf.create(rows, cols, type);
 
681
    return buf;
 
682
}
 
683
 
 
684
#endif
 
685
 
 
686
////////////////////////////////////////////////////////////////
 
687
// Event
 
688
 
 
689
#ifndef HAVE_CUDA
 
690
 
 
691
class cv::cuda::Event::Impl
 
692
{
 
693
public:
 
694
    Impl(unsigned int)
 
695
    {
 
696
        throw_no_cuda();
 
697
    }
 
698
};
 
699
 
 
700
#else
 
701
 
 
702
class cv::cuda::Event::Impl
 
703
{
 
704
public:
 
705
    cudaEvent_t event;
 
706
    bool ownEvent;
 
707
 
 
708
    explicit Impl(unsigned int flags);
 
709
    explicit Impl(cudaEvent_t event);
 
710
    ~Impl();
 
711
};
 
712
 
 
713
cv::cuda::Event::Impl::Impl(unsigned int flags) : event(0), ownEvent(false)
 
714
{
 
715
    cudaSafeCall( cudaEventCreateWithFlags(&event, flags) );
 
716
    ownEvent = true;
 
717
}
 
718
 
 
719
cv::cuda::Event::Impl::Impl(cudaEvent_t e) : event(e), ownEvent(false)
 
720
{
 
721
}
 
722
 
 
723
cv::cuda::Event::Impl::~Impl()
 
724
{
 
725
    if (event && ownEvent)
 
726
    {
 
727
        cudaEventDestroy(event);
 
728
    }
 
729
}
 
730
 
 
731
cudaEvent_t cv::cuda::EventAccessor::getEvent(const Event& event)
 
732
{
 
733
    return event.impl_->event;
 
734
}
 
735
 
 
736
Event cv::cuda::EventAccessor::wrapEvent(cudaEvent_t event)
 
737
{
 
738
    return Event(makePtr<Event::Impl>(event));
 
739
}
 
740
 
 
741
#endif
 
742
 
 
743
cv::cuda::Event::Event(CreateFlags flags)
 
744
{
 
745
#ifndef HAVE_CUDA
 
746
    (void) flags;
 
747
    throw_no_cuda();
 
748
#else
 
749
    impl_ = makePtr<Impl>(flags);
 
750
#endif
 
751
}
 
752
 
 
753
void cv::cuda::Event::record(Stream& stream)
 
754
{
 
755
#ifndef HAVE_CUDA
 
756
    (void) stream;
 
757
    throw_no_cuda();
 
758
#else
 
759
    cudaSafeCall( cudaEventRecord(impl_->event, StreamAccessor::getStream(stream)) );
 
760
#endif
 
761
}
 
762
 
 
763
bool cv::cuda::Event::queryIfComplete() const
 
764
{
 
765
#ifndef HAVE_CUDA
 
766
    throw_no_cuda();
 
767
    return false;
 
768
#else
 
769
    cudaError_t err = cudaEventQuery(impl_->event);
 
770
 
 
771
    if (err == cudaErrorNotReady || err == cudaSuccess)
 
772
        return err == cudaSuccess;
 
773
 
 
774
    cudaSafeCall(err);
 
775
    return false;
 
776
#endif
 
777
}
 
778
 
 
779
void cv::cuda::Event::waitForCompletion()
 
780
{
 
781
#ifndef HAVE_CUDA
 
782
    throw_no_cuda();
 
783
#else
 
784
    cudaSafeCall( cudaEventSynchronize(impl_->event) );
 
785
#endif
 
786
}
 
787
 
 
788
float cv::cuda::Event::elapsedTime(const Event& start, const Event& end)
 
789
{
 
790
#ifndef HAVE_CUDA
 
791
    (void) start;
 
792
    (void) end;
 
793
    throw_no_cuda();
 
794
    return 0.0f;
 
795
#else
 
796
    float ms;
 
797
    cudaSafeCall( cudaEventElapsedTime(&ms, start.impl_->event, end.impl_->event) );
 
798
    return ms;
 
799
#endif
 
800
}