Issue 1:
System information (version)

OpenCV: 3.3

Operating System / Platform: ubuntu16

Compiler: cmake & gcc5.4

Detailed description

opencv knn: dead loop

Steps to reproduce

#include <opencv2/imgproc.hpp>
#include <opencv2/ml.hpp>
#include <vector>
#include "opencv2/core/core_c.h"

using namespace std;
using namespace cv;
using namespace cv::ml;

using cv::Ptr;

int main(){

Mat xTrainData = (Mat_<float>(5,2) << 1, 1, 1, 1, 2, 2, 2, 2, 2, 2);
Mat yTrainLabels = (Mat_<float>(5,1) << 1, 1, 2, 2, 2);

cout << "===============xTrainData===================="<< endl;
cout << "xTrainData: " << xTrainData<< endl;

cout << "===============yTrainLabels===================="<< endl;
cout << "yTrainLabels: " << yTrainLabels << endl;

// KNearest KDTree implementation
Ptr<KNearest> knnKdt = KNearest::create();
knnKdt->setAlgorithmType(KNearest::KDTREE);
knnKdt->setIsClassifier(true);

// When runs to this step, it does not run any more.
knnKdt->train(xTrainData, ml::ROW_SAMPLE, yTrainLabels);

float d_test[2][2] = {{1.3, 1.3}, {2.2, 2.4}};

Mat xTestData = (Mat_<float>(2,1) << 1.1, 1.1, 2, 2.2);
Mat zBestLabels;
cout << "===============xTestData===================="<< endl;
cout << "xTestData: " << xTestData << endl;

knnKdt->findNearest(xTestData, 1, zBestLabels);
cout << "===============zBestLabels===================="<< endl;
cout << "zBestLabels: " << zBestLabels << endl;
}

When I runs it in clion:

What's wrong?

Issue 2:
System information (version)

OpenCV =>4.0.0 beta

Operating System / Platform => Windows 10 64 Bit

Compiler => Visual Studio 2017

Cuda => 10.0

Driver => 416.34

GPU => RTX 2080

Detailed description

The example below crashes when using RTX 2080 but works with GTX 1080.

The example uses the opencv_contrib modules but the bug is in the files scan.hpp and warp_shuffle.hpp in opencv/modules/core/include/opencv2/cuda/

The crash is coming from the kernel hanging inside warpScanInclusive call to cv::cuda::device::shfl_up(iData, i) on line 191 of scan.hpp

On line 56 of warp_shuffle.hpp there is a #define for __shfl_up(x, y, z) to call __shfl_up_sync(0xFFFFFFFFU, x, y, z)

Nvidia must have changed the sync behavior with the new cards and the __shfl_up_sync fails because the mask is being passed to sync all warp threads when all warp threads are not active doe to the check on line 238 of scan.hpp

if (tid < (tiNumScanThreads / OPENCV_CUDA_WARP_SIZE) )// not all threads active inside here
{
//grab top warp elements
T val = s_Data[tid];
//calculate exclusive scan and write back to shared memory
s_Data[tid] = warpScanExclusive(val, s_Data, tid);
}

Nvidia specifically says not to use 0xFFFFFFFF mask when updating legacy code.

https://devblogs.nvidia.com/using-cuda-warp-level-primitives/

I have a workaround for scan.hpp that works for me but I am not sure how to fix it to work more generally in the entire code base.

template <typename T>
__device__ T warpScanInclusive(T idata, volatile T* s_Data, unsigned int tid)
{
#if __CUDA_ARCH__ >= 300
const unsigned int laneId = cv::cuda::device::Warp::laneId();

unsigned mask = __activemask();// need to get mask for active threads

// scan on shuffl functions
#pragma unroll
for (int i = 1; i <= (OPENCV_CUDA_WARP_SIZE / 2); i *= 2)
{
const T n = __shfl_up_sync(mask, idata, i, 32);//need to call with proper mask
if (laneId >= i)
idata += n;
}

return idata;
#else
unsigned int pos = 2 * tid - (tid & (OPENCV_CUDA_WARP_SIZE - 1));
s_Data[pos] = 0;
pos += OPENCV_CUDA_WARP_SIZE;
s_Data[pos] = idata;

s_Data[pos] += s_Data[pos - 1];
s_Data[pos] += s_Data[pos - 2];
s_Data[pos] += s_Data[pos - 4];
s_Data[pos] += s_Data[pos - 8];
s_Data[pos] += s_Data[pos - 16];

return s_Data[pos];
#endif
}

Steps to reproduce

Running the following code runs the cuda kernel that hangs when using a RTX 2080.

#include "opencv2\highgui.hpp"
#include "opencv2\cudaimgproc.hpp"

int main(int argc, char **argv)
{
cv::Ptr<cv::cuda::CLAHE> clahe = cv::cuda::createCLAHE();

const cv::Size size(640, 480);

cv::cuda::GpuMat srcDevice(size, CV_8UC1);
cv::cuda::GpuMat dstDevice(size, CV_8UC1);

cv::cuda::HostMem srcHost(size, CV_8UC1);
cv::cuda::HostMem dstHost(size, CV_8UC1);

cv::randu(srcHost, cv::Scalar::all(0.0), cv::Scalar::all(255.0));
srcDevice.upload(srcHost);

clahe->apply(srcDevice, dstDevice);

dstDevice.download(dstHost);

cv::imshow("src", srcHost);
cv::imshow("dst", dstHost);

cv::waitKey();

return 0;
}