diff --git a/modules/gpu/cuda/Stereo.cu b/modules/gpu/cuda/stereobm.cu
similarity index 95%
rename from modules/gpu/cuda/Stereo.cu
rename to modules/gpu/cuda/stereobm.cu
index a316bb84aa..396248570e 100644
--- a/modules/gpu/cuda/Stereo.cu
+++ b/modules/gpu/cuda/stereobm.cu
@@ -42,10 +42,6 @@
 
 #include "cuda_shared.hpp"
 
-using namespace cv::gpu;
-
-#define cudaSafeCall
-
 #define ROWSperTHREAD 21     // the number of rows a thread will process
 #define BLOCK_W 128          // the thread block width (464)
 #define N_DISPARITIES 8
@@ -218,7 +214,7 @@ __device__ void InitColSSD(int x_tex, int y_tex, int im_pitch, unsigned char* im
     col_ssd[7 * SHARED_MEM_SIZE] = diffa[7];
 }
 
-extern "C" __global__ void stereoKernel(uchar *left, uchar *right, size_t img_step, uchar* disp, size_t disp_pitch, int maxdisp)
+extern "C" __global__ void stereoKernel(unsigned char *left, unsigned char *right, size_t img_step, unsigned char* disp, size_t disp_pitch, int maxdisp)
 {
     extern __shared__ unsigned int col_ssd_cache[];
     unsigned int *col_ssd = col_ssd_cache + BLOCK_W + threadIdx.x;    
@@ -231,7 +227,7 @@ extern "C" __global__ void stereoKernel(uchar *left, uchar *right, size_t img_st
     //int Y = blockIdx.y * ROWSperTHREAD + RADIUS;
 
     unsigned int* minSSDImage = cminSSDImage + X + Y * cminSSD_step;
-    uchar* disparImage = disp + X + Y * disp_pitch;
+    unsigned char* disparImage = disp + X + Y * disp_pitch;
  /*   if (X < cwidth)
     {        
         unsigned int *minSSDImage_end = minSSDImage + min(ROWSperTHREAD, cheight - Y) * minssd_step;
@@ -301,6 +297,7 @@ extern "C" void cv::gpu::impl::stereoBM_GPU(const DevMem2D& left, const DevMem2D
     
     size_t smem_size = (BLOCK_W + N_DISPARITIES * SHARED_MEM_SIZE) * sizeof(unsigned int);      
 
+#define cudaSafeCall
     cudaSafeCall( cudaMemset2D(disp.ptr, disp.step, 0, disp.cols, disp. rows) );
     cudaSafeCall( cudaMemset2D(minSSD_buf.ptr, minSSD_buf.step, 0xFF, minSSD_buf.cols * minSSD_buf.elemSize(), disp. rows) );        
 
diff --git a/modules/gpu/include/opencv2/gpu/gpu.hpp b/modules/gpu/include/opencv2/gpu/gpu.hpp
index 989aa15d3e..4c343014ce 100644
--- a/modules/gpu/include/opencv2/gpu/gpu.hpp
+++ b/modules/gpu/include/opencv2/gpu/gpu.hpp
@@ -56,9 +56,7 @@ namespace cv
         CV_EXPORTS string getDeviceName(int device);
         CV_EXPORTS void setDevice(int device);        
 
-        enum { CV_GPU_CC_10, CV_GPU_CC_11, CV_GPU_CC_12, CV_GPU_CC_13, CV_GPU_CC_20 };
-
-        CV_EXPORTS int getComputeCapability(int device);
+        CV_EXPORTS void getComputeCapability(int device, int* major, int* minor);
         CV_EXPORTS int getNumberOfSMs(int device);
  
         //////////////////////////////// GpuMat ////////////////////////////////
diff --git a/modules/gpu/src/initialization.cpp b/modules/gpu/src/initialization.cpp
index c1c3b9964a..d67d255f5b 100644
--- a/modules/gpu/src/initialization.cpp
+++ b/modules/gpu/src/initialization.cpp
@@ -64,27 +64,15 @@ CV_EXPORTS void cv::gpu::setDevice(int device)
     cudaSafeCall( cudaSetDevice( device ) );
 }
 
-CV_EXPORTS int cv::gpu::getComputeCapability(int device)
-{    
-    cudaDeviceProp prop;
+CV_EXPORTS void cv::gpu::getComputeCapability(int device, int* major, int* minor)
+{
+    cudaDeviceProp prop;    
     cudaSafeCall( cudaGetDeviceProperties( &prop, device) );
 
-    if (prop.major == 2)
-        return CV_GPU_CC_20;
-
-    if (prop.major == 1)
-        switch (prop.minor)
-        {
-        case 0: return CV_GPU_CC_10;
-        case 1: return CV_GPU_CC_11;
-        case 2: return CV_GPU_CC_12;
-        case 3: return CV_GPU_CC_13;        
-        }
-    
-    return -1;
+    *major = prop.major;
+    *minor = prop.minor;
 }
 
-
 CV_EXPORTS int cv::gpu::getNumberOfSMs(int device)
 {
     cudaDeviceProp prop;
diff --git a/modules/gpu/src/precomp.hpp b/modules/gpu/src/precomp.hpp
index e6cfbaa395..deff3df08d 100644
--- a/modules/gpu/src/precomp.hpp
+++ b/modules/gpu/src/precomp.hpp
@@ -55,7 +55,11 @@
 
 #include "cuda_shared.hpp"
 
-#include "cuda_runtime.h"
+#if _MSC_VER >= 1200
+#pragma warning (disable : 4100 4211 4201 4408)
+#endif
+
+#include "cuda_runtime_api.h"
 
 
 #define cudaSafeCall(err)  __cudaSafeCall(err, __FILE__, __LINE__)
diff --git a/modules/gpu/src/stereobm_gpu.cpp b/modules/gpu/src/stereobm_gpu.cpp
index 306e7e651c..a648494bda 100644
--- a/modules/gpu/src/stereobm_gpu.cpp
+++ b/modules/gpu/src/stereobm_gpu.cpp
@@ -49,7 +49,8 @@ using namespace cv::gpu;
 StereoBM_GPU::StereoBM_GPU() : preset(BASIC_PRESET), ndisp(64)  {}
 StereoBM_GPU::StereoBM_GPU(int preset_, int ndisparities_) : preset(preset_), ndisp(ndisparities_) 
 {
-    CV_Assert(ndisp <= std::numeric_limits<unsigned char>::max());
+    const int max_supported_ndisp = 1 << (sizeof(unsigned char) * 8);
+    CV_Assert(ndisp <= max_supported_ndisp);
 }
   
 void StereoBM_GPU::operator() ( const GpuMat& left, const GpuMat& right, GpuMat& disparity) const