CudaNdarray.__setitem now accept an ndarray as a value. It copy the data…

theano/sandbox/cuda/cuda_ndarray.cu

@@ -1429,12 +1429,65 @@ CudaNdarray_Subscript(PyObject * py_self, PyObject * key)
 // See http://docs.python.org/dev/py3k/c-api/object.html#PyObject_SetItem
 // Doesn't handle broadcasting, e.g. a[:] = 5
 // Can only be assigned from a CudaNdarray on the right side
+// Or a ndarray when the left side part is c contiguous.
 static int
 CudaNdarray_setitem(PyObject *o, PyObject  *key, PyObject  *v)
 {
+    if(CudaNdarray_Check(o)  && PyArray_Check(v)){
+        // We try to copy directly into this CudaNdarray from the ndarray
+        CudaNdarray* rval = (CudaNdarray*)CudaNdarray_Subscript(o, key);
+	int typenum = PyArray_TYPE(v);
+	
+	if(!rval){
+            // CudaNdarray_Subscript failed and set the error msg.
+	    Py_XDECREF(rval);
+	    return -1;
+	}
+	if (typenum != REAL_TYPENUM){
+	    PyErr_SetString(PyExc_TypeError, "CudaNdarray.__setitem__: can only copy from float32 arrays");
+	    Py_XDECREF(rval);
+	    return -1;
+	}
+	if(! CudaNdarray_is_c_contiguous(rval)){
+            PyErr_SetString(PyExc_NotImplementedError, "CudaNdarray.__setitem__: When the new value is an ndarray the part where we copy it to must be c contiguous.");
+	    Py_XDECREF(rval);
+	    return -1;
+	}
+	if(rval->nd != ((PyArrayObject*)v)->nd){
+            PyErr_Format(PyExc_NotImplementedError, "CudaNdarray.__setitem__: need same number of dims. destination nd=%d, source nd=%d. No broadcasting implemented.",
+			    rval->nd,((PyArrayObject*)v)->nd);
+	    Py_XDECREF(rval);
+	    return -1;
+	}
+	for(int i=0 ; i<rval->nd ; i++){
+	  if(CudaNdarray_HOST_DIMS(rval)[i] != ((PyArrayObject*)v)->dimensions[i]){
+	    PyErr_Format(PyExc_ValueError, "CudaNdarray.__setitem__: need same dimensions for dim %d, destination=%d, source=%d",i,
+			    CudaNdarray_HOST_DIMS(rval)[i],
+			    ((PyArrayObject*)v)->dimensions[i]);
+	    Py_XDECREF(rval);
+	    return -1;
+	  }
+	}
+	PyArrayObject * py_v = (PyArrayObject*)PyArray_ContiguousFromAny((PyObject*)v, typenum,
+						      rval->nd, rval->nd);
+	cublasSetVector(PyArray_SIZE(py_v),
+			sizeof(real), 
+			PyArray_DATA(py_v), 1,
+			rval->devdata, 1);
+	CNDA_THREAD_SYNC;
+	Py_XDECREF(py_v);
+	Py_XDECREF(rval);
+	if (CUBLAS_STATUS_SUCCESS != cublasGetError()){
+	  PyErr_SetString(PyExc_RuntimeError, "CudaNdarray.__setitem__: error copying ndarray data to device memory");
+	  return -1;
+	}
+	  return 0;
+    }
+
+
     if(!CudaNdarray_Check(o) || !CudaNdarray_Check(v))
     {
-        PyErr_SetString(PyExc_TypeError, "both left and right of setitem must be CudaNdarrays");
+        PyErr_SetString(PyExc_TypeError, "CudaNdarray.__setitem__: left must be a CudaNdarrays and right must be a CudaNdarrays or ndarray");
         return -1;
     }
 
@@ -1458,7 +1511,7 @@ CudaNdarray_setitem(PyObject *o, PyObject  *key, PyObject  *v)
 
     if (cnda_copy_structure_to_device(rval))
     {
-        PyErr_SetString(PyExc_RuntimeError, "CudaNdarray_setitem: syncing structure to device failed");
+        PyErr_SetString(PyExc_RuntimeError, "CudaNdarray.__setitem__: syncing structure to device failed");
         Py_DECREF(rval);
         return -1;
     }

theano/sandbox/cuda/tests/test_cuda_ndarray.py

@@ -387,6 +387,24 @@ def test_gemm_vector_vector():
 
 # ---------------------------------------------------------------------
 
+def test_setitem_matrixscalar0():
+    a = theano._asarray([[0,1,2], [3,4,5]], dtype='float32')
+    _a = cuda_ndarray.CudaNdarray(a)
+
+    b = theano._asarray(8, dtype='float32')
+    _b = cuda_ndarray.CudaNdarray(b)
+
+    # set an element to 8
+    _a[1,1] = _b    
+    a[1,1] = b
+    assert numpy.allclose(a,numpy.asarray(_a))
+
+    #test direct transfert from numpy
+    _a[1,1] = theano._asarray(888, dtype='float32')
+    a[1,1] = theano._asarray(888, dtype='float32')
+    assert numpy.allclose(a,numpy.asarray(_a))
+    
+
 def test_setitem_matrixvector1():
     a = theano._asarray([[0,1,2], [3,4,5]], dtype='float32')
     _a = cuda_ndarray.CudaNdarray(a)
@@ -396,8 +414,22 @@ def test_setitem_matrixvector1():
 
     # set second column to 8,9
     _a[:,1] = _b
-
-    assert numpy.all(numpy.asarray(_a[:,1]) == b)
+    a[:,1] = b
+    assert numpy.allclose(a,numpy.asarray(_a))
+    
+    #test direct transfert from numpy
+    try:
+        _a[:,1] =  b*100
+        a[:,1] =  b*100
+	raise Exception("CudaNdarray.__setitem__ should have returned an error")
+        assert numpy.allclose(a,numpy.asarray(_a))
+    except NotImplementedError, e:
+        pass
+
+    row = theano._asarray([777,888,999], dtype='float32')
+    _a[1,:] = row
+    a[1,:] = row
+    assert numpy.allclose(a,numpy.asarray(_a))
 
 def test_setitem_matrix_tensor3():
     a = numpy.arange(27)
@@ -411,7 +443,25 @@ def test_setitem_matrix_tensor3():
     # set middle row through cube to 7,8,9
     _a[:,1,1] = _b
 
-    assert numpy.all(numpy.asarray(_a[:,1,1]) == b)
+    a[:,1,1] = b
+    assert numpy.allclose(a,numpy.asarray(_a))
+
+    #test direct transfert from numpy
+    try:
+        _a[:,1,1] = b*100
+        a[:,1,1] = b*100
+	raise Exception("CudaNdarray.__setitem__ should have returned an error")  
+        assert numpy.allclose(a,numpy.asarray(_a))
+    except NotImplementedError:
+        pass
+
+    row = theano._asarray([777,888,999], dtype='float32')
+    _a[1,1,:] = row
+    a[1,1,:] = row
+    assert numpy.allclose(a,numpy.asarray(_a))
+
+def test_setitem_from_numpy_error():
+    pass
 
 def test_setitem_matrix_bad_shape():
     a = numpy.arange(27)
@@ -423,13 +473,65 @@ def test_setitem_matrix_bad_shape():
     _b = cuda_ndarray.CudaNdarray(b)
 
     try:
-        # attempt to assign the ndarray b with setitem                                                                                                                                              
-        _a[:,:,1] = _b
+        # attempt to assign the ndarray b with setitem
+        _a[:,1,1] = _b
         assert False
     except ValueError, e:
         #print e
         assert True
 
+    #test direct transfert from numpy
+    try:
+        # attempt to assign the ndarray b with setitem
+        _a[1,1,:] = b
+        assert False
+    except ValueError, e:
+        #print e
+        assert True
+
+def test_setitem_matrix_bad_ndim():
+    a = numpy.arange(27)
+    a.resize((3,3,3))
+    a = theano._asarray(a, dtype='float32')
+    _a = cuda_ndarray.CudaNdarray(a)
+
+    b = theano._asarray([7,8], dtype='float32')
+    _b = cuda_ndarray.CudaNdarray(b)
+
+    try:
+        # attempt to assign the ndarray b with setitem
+        _a[:,:,1] = _b
+        assert False
+    except NotImplementedError, e:
+        #print e
+        assert True
+
+    #test direct transfert from numpy
+    try:
+        # attempt to assign the ndarray b with setitem
+        _a[1,:,:] = b
+        assert False
+    except NotImplementedError, e:
+        #print e
+        assert True
+
+def test_setitem_matrix_bad_type():
+    a = numpy.arange(27)
+    a.resize((3,3,3))
+    a = theano._asarray(a, dtype='float32')
+    _a = cuda_ndarray.CudaNdarray(a)
+
+    b = theano._asarray([7,8], dtype='float64')
+
+    #test direct transfert from numpy
+    try:
+        # attempt to assign the ndarray b with setitem
+        _a[1,:,:] = b
+        assert False
+    except TypeError, e:
+        #print e
+        assert True
+
 def test_setitem_assign_to_slice():
     a = numpy.arange(27)
     a.resize((3,3,3))
@@ -446,7 +548,14 @@ def test_setitem_assign_to_slice():
     # (this corresponds to middle row of matrix _c)
     _c[:,1] = _b
 
-    assert numpy.all(numpy.asarray(_a[:,1,1]) == b)
+    a[:,:,1][:,1] = b
+    assert numpy.allclose(a,numpy.asarray(_a))
+
+    #test direct transfert from numpy
+    _d = _a[1,:,:]
+    _d[1,:] = b*10
+    a[1,:,:][1,:] = b*10
+    assert numpy.allclose(a,numpy.asarray(_a))
 
 def test_setitem_broadcast():
     #test scalar to vector without stride
@@ -485,6 +594,52 @@ def test_setitem_broadcast():
     a[:,:,1] = b.reshape((1,3))
     assert numpy.allclose(numpy.asarray(_a),a)
 
+#This is not supported for now.
+def test_setitem_broadcast_numpy():
+    #test scalar to vector without stride
+    a = numpy.arange(3)
+    a = theano._asarray(a, dtype='float32')
+    _a = cuda_ndarray.CudaNdarray(a)
+
+    b = theano._asarray(9, dtype='float32')
+    try:
+        _a[:] = b.reshape((1,))
+        a[:] = b.reshape((1,))
+	assert False
+        assert numpy.allclose(numpy.asarray(_a),a)
+    except ValueError:
+	    pass
+    #test vector to matrice without stride
+    a = numpy.arange(9)
+    a.resize((3,3))
+    a = theano._asarray(a, dtype='float32')
+    _a = cuda_ndarray.CudaNdarray(a)
+
+    try:
+        b = theano._asarray([7,8,9], dtype='float32')
+        _a[:,:] = b.reshape((1,3))
+        a[:,:] = b.reshape((1,3))
+	assert False
+	assert numpy.allclose(numpy.asarray(_a),a)
+    except ValueError:
+	    pass
+
+    #test vector to matrice with stride
+    a = numpy.arange(27)
+    a.resize((3,3,3))
+    a = theano._asarray(a, dtype='float32')
+    _a = cuda_ndarray.CudaNdarray(a)
+
+    try:
+        b = theano._asarray([[7,8,9],[10,11,12]], dtype='float32')
+        b = b[0]
+        _a[1,:,:] = b.reshape((1,3))
+        a[1,:,:] = b.reshape((1,3))
+	assert False
+	assert numpy.allclose(numpy.asarray(_a),a)
+    except ValueError:
+	    pass
+
 # this also fails for the moment
 def test_setitem_rightvalue_ndarray_fails():
     """
@@ -498,11 +653,20 @@ def test_setitem_rightvalue_ndarray_fails():
     b = theano._asarray([7,8,9], dtype='float32')
     _b = cuda_ndarray.CudaNdarray(b)
 
+    try:
+        # attempt to assign the ndarray b with setitem
+        _a[:,:,1] = _b
+        assert False
+    except NotImplementedError, e:
+        #print e
+        assert True
+
+    #test direct transfert from numpy
     try:
         # attempt to assign the ndarray b with setitem
         _a[:,:,1] = b
         assert False
-    except TypeError, e:
+    except NotImplementedError, e:
         #print e
         assert True