Added warnings about CUBLAS bug to GpuCorrMM documentation

doc/library/tensor/nnet/conv.txt

@@ -82,6 +82,7 @@ TODO: Give examples on how to use these things! They are pretty complicated.
       This is not enabled by default because it uses some extra memory, but the
       overhead is small compared to conv2d_fft, there are no restrictions on
       input or kernel shapes and it is sometimes still faster than cuda-convnet.
+      If using it, please see the warning about a bug in CUDA 5.0 to 6.0 below.
       To enable it for just one Theano function:
 
       .. code-block:: python

theano/sandbox/cuda/blas.py

@@ -785,6 +785,12 @@ class GpuCorrMM(BaseGpuCorrMM):
         faster, but note that it computes a correlation -- if you need to
         compute a convolution, flip the filters as `filters[:,:,::-1,::-1]`.
 
+    :warning: For 700 series Nvidia GPUs of compute capability 3.5 and CUDA 5.0
+        to 6.0, there is a bug in CUBLAS' matrix multiplication function that
+        can make GpuCorrMM or its gradients crash for some input and filter
+        shapes. So if you have a Tesla K20, Tesla K40, Quadro K6000, GeForce GT
+        640 (DDR5), GeForce GTX 780 (or Ti), GeForce GTX TITAN (or Black or Z)
+        and experience a crash, switching to CUDA 6.5 or CUDA 4.2 should fix it.
     """
     def __init__(self, border_mode="valid",
             subsample=(1, 1),

theano/sandbox/cuda/conv_gemm.cu

@@ -254,7 +254,9 @@ CudaNdarray* corrMM(CudaNdarray *const bottom,
             cudaError_t err = cudaGetLastError();
             if (err != cudaSuccess) {
                 PyErr_Format(PyExc_RuntimeError,
-                             "GpuCorrMM encountered a CUDA error in im2col: %s\n",
+                             "GpuCorrMM encountered a CUDA error in im2col: %s\n"
+                             "This could be a known bug in CUDA, please see the "
+                             "GpuCorrMM() documentation.\n",
                              cudaGetErrorString(err));
                 return NULL;
             }
@@ -269,7 +271,9 @@ CudaNdarray* corrMM(CudaNdarray *const bottom,
                     top->devdata + n * top_stride, N_);
             if (status != CUBLAS_STATUS_SUCCESS) {
                 PyErr_Format(PyExc_RuntimeError,
-                        "GpuCorrMM encountered a CUBLAS error: %s\n",
+                        "GpuCorrMM encountered a CUBLAS error: %s\n"
+                        "This could be a known bug in CUDA, please see the "
+                        "GpuCorrMM() documentation.\n",
                         cublasGetErrorString(status));
                 return NULL;
             }
@@ -313,7 +317,9 @@ CudaNdarray* corrMM(CudaNdarray *const bottom,
                 sizeof(float) * M_ * K_);
         if (err != cudaSuccess) {
                 PyErr_Format(PyExc_RuntimeError,
-                        "GpuCorrMM encountered a CUDA error in cudaMemsetAsync: %s\n",
+                        "GpuCorrMM encountered a CUDA error in cudaMemsetAsync: %s\n"
+                        "This could be a known bug in CUDA, please see the "
+                        "GpuCorrMM() documentation.\n",
                         cudaGetErrorString(err));
                 return NULL;
         }
@@ -325,7 +331,9 @@ CudaNdarray* corrMM(CudaNdarray *const bottom,
             err = cudaGetLastError();
             if (err != cudaSuccess) {
                 PyErr_Format(PyExc_RuntimeError,
-                             "GpuCorrMM encountered a CUDA error in im2col: %s\n",
+                             "GpuCorrMM encountered a CUDA error in im2col: %s\n"
+                             "This could be a known bug in CUDA, please see the "
+                             "GpuCorrMM() documentation.\n",
                              cudaGetErrorString(err));
                 return NULL;
             }
@@ -340,7 +348,9 @@ CudaNdarray* corrMM(CudaNdarray *const bottom,
                     weight->devdata, K_);
             if (status != CUBLAS_STATUS_SUCCESS) {
                 PyErr_Format(PyExc_RuntimeError,
-                        "GpuCorrMM encountered a CUBLAS error: %s\n",
+                        "GpuCorrMM encountered a CUBLAS error: %s\n"
+                        "This could be a known bug in CUDA, please see the "
+                        "GpuCorrMM() documentation.\n",
                         cublasGetErrorString(status));
                 return NULL;
             }
@@ -388,7 +398,9 @@ CudaNdarray* corrMM(CudaNdarray *const bottom,
                     col->devdata, N_);
             if (status != CUBLAS_STATUS_SUCCESS) {
                 PyErr_Format(PyExc_RuntimeError,
-                        "GpuCorrMM encountered a CUBLAS error: %s\n",
+                        "GpuCorrMM encountered a CUBLAS error: %s\n"
+                        "This could be a known bug in CUDA, please see the "
+                        "GpuCorrMM() documentation.\n",
                         cublasGetErrorString(status));
                 return NULL;
             }
@@ -398,7 +410,9 @@ CudaNdarray* corrMM(CudaNdarray *const bottom,
             cudaError_t err = cudaGetLastError();
             if (err != cudaSuccess) {
                 PyErr_Format(PyExc_RuntimeError,
-                             "GpuCorrMM encountered a CUDA error in col2im: %s\n",
+                             "GpuCorrMM encountered a CUDA error in col2im: %s\n"
+                             "This could be a known bug in CUDA, please see the "
+                             "GpuCorrMM() documentation.\n",
                              cudaGetErrorString(err));
                 return NULL;
             }