<h5 dir="auto">System information (version)</h5>
<ul dir="auto">
<li>OpenCV: 3.3</li>
<li>Operating System / Platform: ubuntu16</li>
<li>Compiler: cmake &amp; gcc5.4</li>
</ul>
<h5 dir="auto">Detailed description</h5>
<p dir="auto">opencv knn: dead loop</p>
<h5 dir="auto">Steps to reproduce</h5>
<div class="highlight highlight-source-c++ notranslate position-relative overflow-auto" dir="auto" data-snippet-clipboard-copy-content="#include &lt;opencv2/imgproc.hpp&gt;
#include &lt;opencv2/ml.hpp&gt;
#include &lt;vector&gt;
#include &quot;opencv2/core/core_c.h&quot;

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

using cv::Ptr;

int main(){

    Mat xTrainData = (Mat_&lt;float&gt;(5,2) &lt;&lt; 1, 1, 1, 1, 2, 2, 2, 2, 2, 2);
    Mat yTrainLabels = (Mat_&lt;float&gt;(5,1) &lt;&lt; 1, 1, 2, 2, 2);

    cout &lt;&lt; &quot;===============xTrainData====================&quot;&lt;&lt; endl;
    cout &lt;&lt; &quot;xTrainData: &quot;  &lt;&lt; xTrainData&lt;&lt; endl;

    cout &lt;&lt; &quot;===============yTrainLabels====================&quot;&lt;&lt; endl;
    cout &lt;&lt; &quot;yTrainLabels: &quot;  &lt;&lt; yTrainLabels &lt;&lt; endl;

    // KNearest KDTree implementation
    Ptr&lt;KNearest&gt; knnKdt = KNearest::create();
    knnKdt-&gt;setAlgorithmType(KNearest::KDTREE);
    knnKdt-&gt;setIsClassifier(true);

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

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

    Mat xTestData = (Mat_&lt;float&gt;(2,1) &lt;&lt; 1.1, 1.1, 2, 2.2);
    Mat zBestLabels;
    cout &lt;&lt; &quot;===============xTestData====================&quot;&lt;&lt; endl;
    cout &lt;&lt; &quot;xTestData: &quot;  &lt;&lt; xTestData &lt;&lt; endl;

    knnKdt-&gt;findNearest(xTestData, 1, zBestLabels);
    cout &lt;&lt; &quot;===============zBestLabels====================&quot;&lt;&lt; endl;
    cout &lt;&lt; &quot;zBestLabels: &quot;  &lt;&lt; zBestLabels &lt;&lt; endl;
}
"><pre class="notranslate">#<span class="pl-k">include</span> <span class="pl-s"><span class="pl-pds">&lt;</span>opencv2/imgproc.hpp<span class="pl-pds">&gt;</span></span>
#<span class="pl-k">include</span> <span class="pl-s"><span class="pl-pds">&lt;</span>opencv2/ml.hpp<span class="pl-pds">&gt;</span></span>
#<span class="pl-k">include</span> <span class="pl-s"><span class="pl-pds">&lt;</span>vector<span class="pl-pds">&gt;</span></span>
#<span class="pl-k">include</span> <span class="pl-s"><span class="pl-pds">"</span>opencv2/core/core_c.h<span class="pl-pds">"</span></span>

<span class="pl-k">using</span> <span class="pl-k">namespace</span> <span class="pl-en">std</span><span class="pl-k">;</span>
<span class="pl-k">using</span> <span class="pl-k">namespace</span> <span class="pl-en">cv</span><span class="pl-k">;</span>
<span class="pl-k">using</span> <span class="pl-k">namespace</span> <span class="pl-en">cv</span>::ml<span class="pl-k">;</span>

<span class="pl-k">using</span> cv::<span class="pl-c1">Ptr</span>;

<span class="pl-k">int</span> <span class="pl-en">main</span>(){

    Mat xTrainData = (Mat_&lt;<span class="pl-k">float</span>&gt;(<span class="pl-c1">5</span>,<span class="pl-c1">2</span>) &lt;&lt; <span class="pl-c1">1</span>, <span class="pl-c1">1</span>, <span class="pl-c1">1</span>, <span class="pl-c1">1</span>, <span class="pl-c1">2</span>, <span class="pl-c1">2</span>, <span class="pl-c1">2</span>, <span class="pl-c1">2</span>, <span class="pl-c1">2</span>, <span class="pl-c1">2</span>);
    Mat yTrainLabels = (Mat_&lt;<span class="pl-k">float</span>&gt;(<span class="pl-c1">5</span>,<span class="pl-c1">1</span>) &lt;&lt; <span class="pl-c1">1</span>, <span class="pl-c1">1</span>, <span class="pl-c1">2</span>, <span class="pl-c1">2</span>, <span class="pl-c1">2</span>);

    cout &lt;&lt; <span class="pl-s"><span class="pl-pds">"</span>===============xTrainData====================<span class="pl-pds">"</span></span>&lt;&lt; endl;
    cout &lt;&lt; <span class="pl-s"><span class="pl-pds">"</span>xTrainData: <span class="pl-pds">"</span></span>  &lt;&lt; xTrainData&lt;&lt; endl;

    cout &lt;&lt; <span class="pl-s"><span class="pl-pds">"</span>===============yTrainLabels====================<span class="pl-pds">"</span></span>&lt;&lt; endl;
    cout &lt;&lt; <span class="pl-s"><span class="pl-pds">"</span>yTrainLabels: <span class="pl-pds">"</span></span>  &lt;&lt; yTrainLabels &lt;&lt; endl;

    <span class="pl-c"><span class="pl-c">//</span> KNearest KDTree implementation</span>
    <span class="pl-c1">Ptr</span>&lt;KNearest&gt; knnKdt = <span class="pl-c1">KNearest::create</span>();
    knnKdt-&gt;<span class="pl-c1">setAlgorithmType</span>(KNearest::KDTREE);
    knnKdt-&gt;<span class="pl-c1">setIsClassifier</span>(<span class="pl-c1">true</span>);

   <span class="pl-c"><span class="pl-c">//</span> When runs to this step, it does not run any more.</span>
    knnKdt-&gt;<span class="pl-c1">train</span>(xTrainData, ml::ROW_SAMPLE, yTrainLabels);

    <span class="pl-k">float</span> d_test[<span class="pl-c1">2</span>][<span class="pl-c1">2</span>] = {{<span class="pl-c1">1.3</span>, <span class="pl-c1">1.3</span>}, {<span class="pl-c1">2.2</span>, <span class="pl-c1">2.4</span>}};

    Mat xTestData = (Mat_&lt;<span class="pl-k">float</span>&gt;(<span class="pl-c1">2</span>,<span class="pl-c1">1</span>) &lt;&lt; <span class="pl-c1">1.1</span>, <span class="pl-c1">1.1</span>, <span class="pl-c1">2</span>, <span class="pl-c1">2.2</span>);
    Mat zBestLabels;
    cout &lt;&lt; <span class="pl-s"><span class="pl-pds">"</span>===============xTestData====================<span class="pl-pds">"</span></span>&lt;&lt; endl;
    cout &lt;&lt; <span class="pl-s"><span class="pl-pds">"</span>xTestData: <span class="pl-pds">"</span></span>  &lt;&lt; xTestData &lt;&lt; endl;

    knnKdt-&gt;<span class="pl-c1">findNearest</span>(xTestData, <span class="pl-c1">1</span>, zBestLabels);
    cout &lt;&lt; <span class="pl-s"><span class="pl-pds">"</span>===============zBestLabels====================<span class="pl-pds">"</span></span>&lt;&lt; endl;
    cout &lt;&lt; <span class="pl-s"><span class="pl-pds">"</span>zBestLabels: <span class="pl-pds">"</span></span>  &lt;&lt; zBestLabels &lt;&lt; endl;
}
</pre></div>
<p dir="auto">When I runs it in clion:</p>
<p dir="auto"><a target="_blank" rel="noopener noreferrer nofollow" href="https://user-images.githubusercontent.com/16144935/42207391-61e4a88a-7edc-11e8-8eb9-7338ae273342.png"><img src="https://user-images.githubusercontent.com/16144935/42207391-61e4a88a-7edc-11e8-8eb9-7338ae273342.png" alt="image" style="max-width: 100%;"></a></p>
<p dir="auto">What's wrong?</p>
<h5 dir="auto">System information (version)</h5>

<ul dir="auto">
<li>OpenCV =&gt;4.0.0 beta</li>
<li>Operating System / Platform =&gt; Windows 10 64 Bit</li>
<li>Compiler =&gt; Visual Studio 2017</li>
<li>Cuda =&gt; 10.0</li>
<li>Driver =&gt; 416.34</li>
<li>GPU =&gt; RTX 2080</li>
</ul>
<h5 dir="auto">Detailed description</h5>

<p dir="auto">The example below crashes when using RTX 2080 but works with GTX 1080.</p>
<p dir="auto">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/</p>
<p dir="auto">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</p>
<p dir="auto">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)</p>
<p dir="auto">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</p>
<div class="highlight highlight-source-c++ notranslate position-relative overflow-auto" dir="auto" data-snippet-clipboard-copy-content="
            if (tid &lt; (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);
            }
"><pre class="notranslate">            <span class="pl-k">if</span> (tid &lt; (tiNumScanThreads / OPENCV_CUDA_WARP_SIZE) )<span class="pl-c"><span class="pl-c">//</span> not all threads active inside here</span>
            {
                <span class="pl-c"><span class="pl-c">//</span>grab top warp elements</span>
                T val = s_Data[tid];
                <span class="pl-c"><span class="pl-c">//</span>calculate exclusive scan and write back to shared memory</span>
                s_Data[tid] = <span class="pl-c1">warpScanExclusive</span>(val, s_Data, tid);
            }
</pre></div>
<p dir="auto">Nvidia specifically says not to use 0xFFFFFFFF mask when updating legacy code.<br>
<a href="https://devblogs.nvidia.com/using-cuda-warp-level-primitives/" rel="nofollow">https://devblogs.nvidia.com/using-cuda-warp-level-primitives/</a></p>
<p dir="auto">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.</p>
<div class="highlight highlight-source-c++ notranslate position-relative overflow-auto" dir="auto" data-snippet-clipboard-copy-content="
    template &lt;typename T&gt;
    __device__ T warpScanInclusive(T idata, volatile T* s_Data, unsigned int tid)
    {
    #if __CUDA_ARCH__ &gt;= 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 &lt;= (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 &gt;= i)
                  idata += n;
        }

        return idata;
    #else
        unsigned int pos = 2 * tid - (tid &amp; (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
    }
"><pre class="notranslate">    <span class="pl-k">template </span>&lt;<span class="pl-k">typename</span> T&gt;
    __device__ T <span class="pl-en">warpScanInclusive</span>(T idata, <span class="pl-k">volatile</span> T* s_Data, <span class="pl-k">unsigned</span> <span class="pl-k">int</span> tid)
    {
    #<span class="pl-k">if</span> __CUDA_ARCH__ &gt;= 300
        <span class="pl-k">const</span> <span class="pl-k">unsigned</span> <span class="pl-k">int</span> laneId = <span class="pl-c1">cv::cuda::device::Warp::laneId</span>();

        <span class="pl-k">unsigned</span> mask = <span class="pl-c1">__activemask</span>();<span class="pl-c"><span class="pl-c">//</span> need to get mask for active threads</span>

        <span class="pl-c"><span class="pl-c">//</span> scan on shuffl functions</span>
        #<span class="pl-k">pragma</span> unroll
        <span class="pl-k">for</span> (<span class="pl-k">int</span> i = <span class="pl-c1">1</span>; i &lt;= (OPENCV_CUDA_WARP_SIZE / <span class="pl-c1">2</span>); i *= <span class="pl-c1">2</span>)
        {
            <span class="pl-k">const</span> T n = <span class="pl-c1">__shfl_up_sync</span>(mask, idata, i, <span class="pl-c1">32</span>);<span class="pl-c"><span class="pl-c">//</span>need to call with proper mask</span>
            <span class="pl-k">if</span> (laneId &gt;= i)
                  idata += n;
        }

        <span class="pl-k">return</span> idata;
    #<span class="pl-k">else</span>
        <span class="pl-k">unsigned</span> <span class="pl-k">int</span> pos = <span class="pl-c1">2</span> * tid - (tid &amp; (OPENCV_CUDA_WARP_SIZE - <span class="pl-c1">1</span>));
        s_Data[pos] = <span class="pl-c1">0</span>;
        pos += OPENCV_CUDA_WARP_SIZE;
        s_Data[pos] = idata;

        s_Data[pos] += s_Data[pos - <span class="pl-c1">1</span>];
        s_Data[pos] += s_Data[pos - <span class="pl-c1">2</span>];
        s_Data[pos] += s_Data[pos - <span class="pl-c1">4</span>];
        s_Data[pos] += s_Data[pos - <span class="pl-c1">8</span>];
        s_Data[pos] += s_Data[pos - <span class="pl-c1">16</span>];

        <span class="pl-k">return</span> s_Data[pos];
    #<span class="pl-k">endif</span>
    }
</pre></div>
<h5 dir="auto">Steps to reproduce</h5>

<p dir="auto">Running the following code runs the cuda kernel that hangs when using a RTX 2080.</p>
<div class="highlight highlight-source-c++ notranslate position-relative overflow-auto" dir="auto" data-snippet-clipboard-copy-content="
#include &quot;opencv2\highgui.hpp&quot;
#include &quot;opencv2\cudaimgproc.hpp&quot;

int main(int argc, char **argv)
{
   cv::Ptr&lt;cv::cuda::CLAHE&gt; 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-&gt;apply(srcDevice, dstDevice);

   dstDevice.download(dstHost);

   cv::imshow(&quot;src&quot;, srcHost);
   cv::imshow(&quot;dst&quot;, dstHost);

   cv::waitKey();

   return 0;
}
"><pre class="notranslate">#<span class="pl-k">include</span> <span class="pl-s"><span class="pl-pds">"</span>opencv2\highgui.hpp<span class="pl-pds">"</span></span>
#<span class="pl-k">include</span> <span class="pl-s"><span class="pl-pds">"</span>opencv2\cudaimgproc.hpp<span class="pl-pds">"</span></span>

<span class="pl-k">int</span> <span class="pl-en">main</span>(<span class="pl-k">int</span> argc, <span class="pl-k">char</span> **argv)
{
   cv::<span class="pl-c1">Ptr</span>&lt;cv::cuda::CLAHE&gt; clahe = <span class="pl-c1">cv::cuda::createCLAHE</span>();

   <span class="pl-k">const</span> cv::<span class="pl-c1">Size</span> <span class="pl-smi">size</span>(<span class="pl-c1">640</span>, <span class="pl-c1">480</span>);

   cv::cuda::GpuMat <span class="pl-smi">srcDevice</span>(size, CV_8UC1);
   cv::cuda::GpuMat <span class="pl-smi">dstDevice</span>(size, CV_8UC1);

   cv::cuda::HostMem <span class="pl-smi">srcHost</span>(size, CV_8UC1);
   cv::cuda::HostMem <span class="pl-smi">dstHost</span>(size, CV_8UC1);

   <span class="pl-c1">cv::randu</span>(srcHost, <span class="pl-c1">cv::Scalar::all</span>(<span class="pl-c1">0.0</span>), <span class="pl-c1">cv::Scalar::all</span>(<span class="pl-c1">255.0</span>));
   srcDevice.<span class="pl-c1">upload</span>(srcHost);

   clahe-&gt;<span class="pl-c1">apply</span>(srcDevice, dstDevice);

   dstDevice.<span class="pl-c1">download</span>(dstHost);

   <span class="pl-c1">cv::imshow</span>(<span class="pl-s"><span class="pl-pds">"</span>src<span class="pl-pds">"</span></span>, srcHost);
   <span class="pl-c1">cv::imshow</span>(<span class="pl-s"><span class="pl-pds">"</span>dst<span class="pl-pds">"</span></span>, dstHost);

   <span class="pl-c1">cv::waitKey</span>();

   <span class="pl-k">return</span> <span class="pl-c1">0</span>;
}
</pre></div>