CUDA Fast detector much slower than normal FAST

Kansai · May 26, 2021, 1:08am

I am using OpenCV FAST Feature detector in both CPU and CUDA versions as follows:

For the normal CPU version

cv::FAST(img,kps, threshold, true);

For the CUDA version

ptrDetector= cv::cuda::FastFeatureDetector::create(threshold, true);

and then later

ptrDetector->setThreshold(threshold);
ptrDetector->detect(img,kps);

I measure how long does it take to process and in my Ubuntu PC it takes

CPU version: 0.424199 ms

CUDA version: 0.871586 ms

I wonder why is this like this? Is CUDA not adequate for FAST detectors?

EDIT:
I just want to point out that there is a similar question to this , but in there the answer was basically “it takes time to upload and download data”

which in my case, is not the answer, because I am timing only the processing time (excluding the upload and download time)

kbarni · May 26, 2021, 8:13am

Take a look at this question too: https://forum.opencv.org/t/opencv-cuda-extremely-slow

Basically, the GPU startup takes some time. For me, the processing of the first image took 20-30x more time as the next images.

I suggest to test several images in a loop and measure the processing time for each image, and see if there’s an improvement.

Kansai · May 26, 2021, 8:27am

Thanks for the reply. I am aware of the time that initialization takes (it was something that surprised me a lot actually)
My current question though refers to the amount of time FAST takes for a lot of images. In my current program I do FAST on 1000 images. CUDA keeps being slow

I tried also on a Jetson AGX Xavier and I got

CPU version 0.89ms
CUDA version 4.51 ms (+0.39ms for upload)

cudawarped · May 26, 2021, 9:41am

In that post I mentioned that it could be the hardware which is limiting you. The spreadsheet I linked to there shows the performance of FAST on lower end GPU’s is not that much faster than on the CPU, although I would still expect a speed increase. It also looks to be image dependent.

Just running the numbers a Jetson AGX Xavier is only marginally faster than a Geforce GTX 650 Ti (NVIDIA Jetson AGX Xavier 16 GB Specs | TechPowerUp GPU Database), a mid range graphics card released in 2012. I can’t comment on the performance of the ARM processor but my guess would be that it is not unthinkable that for this algorithm the performance would be superior.

Additionally did you read the post @kbarni linked to where we discussed other factors which could be limiting performance.

I guess if everything is fully optimized in your set up then the answer would be no on your hardware but yes on faster GPU’s. For example a state of the art GPU could have ~25x the performance if the size of the images and/or the implementation was sufficient to saturate the GPU.

Kansai · May 26, 2021, 9:56am

Thanks. I am in the process of reading the post. Some of the suggestions (not measuring upload and download for example) are originally done, so doesn’t help much but there is one (pre-allocating) that I want to try. Oh, and I am using C++.

The phenomena that CUDA is slower happens both in the Xavier and in my Ubuntu Host that has a GeForce RTX 2080 Mobile

cudawarped · May 26, 2021, 10:00am

Can you send me a code snipet, I have the same GPU?

Kansai · May 26, 2021, 10:06am

Can you send me a code snipet, I have the same GPU?

I will simplify the code a bit so that it takes only one cpp file and send it tomorrow thanks

cudawarped · May 27, 2021, 1:29pm

Thank you for sending the code over I have had a quick look on my machine to see what times I get. I made two slight alterations to your code:

I averaged 100 iterations because the CPU hi res timers were slightly pessimistic which I confirmed using event timers on the GPU.
I removed std::cout from the timing as this CPU operation should not be included.

With the parameters you are using on a 1500 x 1500 image of a checkerboard the execution time for both the CPU and GPU was almost the same with the CPU slightly faster. i7-8700 vs RTX 2080 mobile.
CPU: 0.84 ms
GPU: 1.00 ms

If I use detectAsync() instead of detect() which also downloads the keypoints and processes them on the CPU then the GPU time is reduced.
GPU: 0.90 ms

I then had a look at the execution time in the Nvidia Visual Profiler and noticed that the execution time is ~0.1 ms. The 0.9ms is loss caused by:

The routine allocates memory on the device and copies the number of keypoints from the device to the host which is included in what you are timing. This requires synchronizing the device with the host and copying from device to host memory which stalls everything, taking longer than the kernel itself.
Including non max suppression. This requires an extra kernel call with extra synchronization and memory overhead.

This is a good example demonstrating that the implementation of the algorithm is as important as the raw performance of the GPU.

If I remove max suppression on both the CPU and GPU the GPU time falls significantly.
CPU 0.71 ms
GPU 0.17 ms

If I use a lower threshold say 20 (which is the number chosen in the inbuilt performance test) then the GPU time does not appear to be affected however the CPU time doubles.
CPU 1.42 ms
GPU 0.14 ms

I know this probably doesn’t help that much but it confirms to me the results from the performance test where the GPU was several times faster than the CPU.

Kansai · May 28, 2021, 12:04am

Thank you very much for your analysis. Today I will go and try to reproduce your experiments. I haven’t use the Nvidia Visual Profiler yet but I will google it. If you have a recommended resource for its use, that would also be very much welcomed.

Just one question, when you said that the profiler notified 0.1ms does that mean that the feature extraction itself only took that time and the rest (0.8ms) was allocation copies and synchronization?

Perhaps tomorrow I will have some results to reflect on. Thanks again

cudawarped · May 28, 2021, 9:08am

The Visual Profiler for CUDA 11.3 can be found

C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v11.3\libnvvp\nvvp.exe

You need to install the JRE for it to work and I you may have to add

C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v11.3\extras\CUPTI\lib64

to your system/user path.
It shows a break down of all the CUDA kernel launches memory operation synchronization etc.
I would look at that break down with reference to

github.com

opencv/opencv_contrib/blob/029c2833623ae658fb1820d5da30e5023134eaf5/modules/cudafeatures2d/src/fast.cpp#L115

    
      
                  return;
              }
          
          
    BufferPool pool(Stream::Null());
              GpuMat d_keypoints = pool.getBuffer(ROWS_COUNT, max_npoints_, CV_32FC1);
          
          
    detectAsync(_image, d_keypoints, _mask, Stream::Null());
              convert(d_keypoints, keypoints);
          }
          
          
void FAST_Impl::detectAsync(InputArray _image, OutputArray _keypoints, InputArray _mask, Stream& stream)
          {
              using namespace cv::cuda::device::fast;
          
          
    cudaSafeCall( cudaMalloc(&d_counter, sizeof(unsigned int)) );
          
          
    const GpuMat img = _image.getGpuMat();
              const GpuMat mask = _mask.getGpuMat();
          
          
    CV_Assert( img.type() == CV_8UC1 );
              CV_Assert( mask.empty() || (mask.type() == CV_8UC1 && mask.size() == img.size()) );

and

github.com

opencv/opencv_contrib/blob/029c2833623ae658fb1820d5da30e5023134eaf5/modules/cudafeatures2d/src/cuda/fast.cu#L282

    
      
                      if (ind < maxKeypoints)
                          kpLoc[ind] = make_short2(j, i);
                  }
              }
          
          
    #endif
          }
          
          
int calcKeypoints_gpu(PtrStepSzb img, PtrStepSzb mask, short2* kpLoc, int maxKeypoints, PtrStepSzi score, int threshold, unsigned int* d_counter, cudaStream_t stream)
          {
              dim3 block(32, 8);
          
          
    dim3 grid;
              grid.x = divUp(img.cols - 6, block.x);
              grid.y = divUp(img.rows - 6, block.y);
          
          
    cudaSafeCall( cudaMemsetAsync(d_counter, 0, sizeof(unsigned int), stream) );
          
          
    if (score.data)
              {
                  if (mask.data)

where you can see calls to cudaMalloc(), cudaFree() and several calls to both cudaMemsetAsync() and cudaStreamSynchronize().

Topic		Replies	Views
OpenCV CUDA extremely slow cuda	3	6279	April 30, 2021
OpenCV Optical Flow Cuda Naiva Implementation Slower then CPU Python cuda	3	313	April 4, 2024
CUDA Fast detector sometimes finds less keypoints from a certain Y features2d , cuda	2	329	June 1, 2021
Some opencv cudafilter functions is slower than CPU code on Jetson Xavier NX C++ filter , cuda , cudaarithm	1	274	November 8, 2023
OpenCV FAST operation. Why use FAST instead of FastFeatureDetector	1	315	May 20, 2021

CUDA Fast detector much slower than normal FAST

Related Topics