Add tests for the grad of the op version (which is broken for now).

theano/sandbox/cuda/blocksparse.py

@@ -50,13 +50,17 @@ def gemv(alpha, A, x, beta, y):
 
 
 def ger(alpha, x, y, A):
-    assert A.shape[0] == x.shape[0]
-    assert A.shape[1] == y.shape[0]
+    print A.shape, x.shape, y.shape
+    assert A.shape[1] == x.shape[0]
+    assert A.shape[0] == y.shape[0]
 
     handle = scikits.cuda.misc._global_cublas_handle
 
-    cublas.cublasSger(handle, A.shape[0], A.shape[1], x.gpudata, x.strides[0],
-                      y.gpudata, y.strides[0], A.gpudata, A.strides[1])
+    cublas.cublasSger(handle, A.shape[0], A.shape[1], alpha,
+                      x.gpudata, x.strides[0],
+                      y.gpudata, y.strides[0],
+                      A.gpudata, A.strides[0])
+
 
 def bptr(a):
     assert (a.ndim == 3 and a.strides[2] == 1)
@@ -109,7 +113,7 @@ class SparseBlockGemvDS(GpuOp):
         go = grads[0]
 
         # might revise that interface to not have a huge output
-        Wgrad = sparse_block_outer_ss(W.zeros_like(), go, h.T,
+        Wgrad = sparse_block_outer_ss(W.zeros_like(), go, h,
                                       outputIdx, inputIdx)
         hgrad = sparse_block_gemv_ds(h.zeros_like(), W.T, go,
                                      outputIdx, inputIdx)
@@ -128,6 +132,9 @@ class SparseBlockOuterSS(GpuOp):
         self.inplace = False
 
     def make_node(self, o, x, y, xIdx, yIdx):
+        o = basic_ops.as_cuda_ndarray_variable(o)
+        x = basic_ops.as_cuda_ndarray_variable(x)
+        y = basic_ops.as_cuda_ndarray_variable(y)
         return Apply(self, [o, x, y, xIdx, yIdx],
                      [o.type()])
 
@@ -142,8 +149,8 @@ class SparseBlockOuterSS(GpuOp):
             out_id = xIdx[i]
             for j in range(y.shape[0]):
                 inp_id = yIdx[j]
-                ger(np.float32(1.0), x[i],
-                    y[j], np.float32(1.0), o[i, j])
+                ger(numpy.float32(1.0), x[i],
+                    y[j], o[inp_id, out_id])
 
         out[0] = o
 

theano/sandbox/cuda/tests/test_blocksparse.py

@@ -64,7 +64,6 @@ def test_blocksparse_op():
     iIdx = tensor.lvector()
     oIdx = tensor.lvector()
 
-
     o = sparse_block_gemv_ds(b.take(oIdx, axis=0), W, h, iIdx, oIdx)
 
     f = theano.function([W, h, iIdx, b, oIdx], o)
@@ -75,3 +74,19 @@ def test_blocksparse_op():
     ref_out = blocksparse(W_val, h_val, iIdx_val, b_val, oIdx_val)
 
     utt.assert_allclose(ref_out, th_out)
+
+
+def test_blocksparse_op_grad():
+    h_val = randn(2, 3).astype('float32')
+    iIdx_val = numpy.random.permutation(3)[:2]
+    oIdx_val = numpy.random.permutation(3)[:2]
+    W_val = randn(3, 3, 3, 3).astype('float32')
+    b_val = randn(3, 3).astype('float32')
+
+    iIdx = theano.tensor.constant(iIdx_val)
+    oIdx = theano.tensor.constant(oIdx_val)
+
+    def f(b, W, h):
+        return sparse_block_gemv_ds(b.take(oIdx, axis=0), W, h, iIdx, oIdx)
+
+    utt.verify_grad(f, [b_val, W_val, h_val])