Fix direct uses of ops.

theano/gpuarray/blockgemv.c

@@ -27,8 +27,7 @@ int APPLY_SPECIFIC(blockgemv)(PyGpuArrayObject *o, PyGpuArrayObject *W,
   gpuarray_blas_ops *blas_ops;
   int err;
 
-  err = ctx->ops->property(ctx->ctx, NULL, NULL,
-                           GA_CTX_PROP_BLAS_OPS, &blas_ops);
+  err = gpucontext_property(ctx->ctx, GA_CTX_PROP_BLAS_OPS, &blas_ops);
   if (err != GA_NO_ERROR) {
     PyErr_SetString(PyExc_RuntimeError, "Can't get blas ops");
     return -1;

theano/gpuarray/blockger.c

@@ -15,8 +15,7 @@ int APPLY_SPECIFIC(blockger)(PyGpuArrayObject *o, PyGpuArrayObject *x,
   gpuarray_blas_ops *blas_ops;
   int err;
 
-  err = ctx->ops->property(ctx->ctx, NULL, NULL,
-                           GA_CTX_PROP_BLAS_OPS, &blas_ops);
+  err = gpucontext_property(ctx->ctx, GA_CTX_PROP_BLAS_OPS, &blas_ops);
   if (err != GA_NO_ERROR) {
     PyErr_SetString(PyExc_RuntimeError, "Can't get blas ops");
     return -1;

theano/gpuarray/dnn_fwd.c

@@ -105,7 +105,7 @@ APPLY_SPECIFIC(conv_fwd)(PyGpuArrayObject *input, PyGpuArrayObject *kerns,
     algo = choice.algo;
 #else
     size_t free;
-    int err2 = c->ops->property(c->ctx, NULL, NULL, GA_CTX_PROP_FREE_GMEM, &free);
+    int err2 = gpucontext_property(c->ctx, GA_CTX_PROP_FREE_GMEM, &free);
 
     if (err2 != GA_NO_ERROR) {
       PyErr_Format(PyExc_RuntimeError, "Error when trying to find the "
@@ -234,7 +234,7 @@ APPLY_SPECIFIC(conv_fwd)(PyGpuArrayObject *input, PyGpuArrayObject *kerns,
      * to place a nice get_work_mem() function in.
      */
     if (worksize != 0) {
-      workspace = c->ops->buffer_alloc(c->ctx, worksize, NULL, 0, NULL);
+      workspace = gpudata_alloc(c->ctx, worksize, NULL, 0, NULL);
       if (workspace == NULL) {
         PyErr_SetString(PyExc_RuntimeError,
                         "Could not allocate working memory");
@@ -258,7 +258,7 @@ APPLY_SPECIFIC(conv_fwd)(PyGpuArrayObject *input, PyGpuArrayObject *kerns,
       APPLY_SPECIFIC(output), PyGpuArray_DEV_DATA(*output));
 
     if (worksize != 0)
-      c->ops->buffer_release(workspace);
+      gpudata_release(workspace);
 
     cuda_record(input->ga.data, GPUARRAY_CUDA_WAIT_READ);
     cuda_record(kerns->ga.data, GPUARRAY_CUDA_WAIT_READ);

theano/gpuarray/dnn_gi.c

@@ -106,7 +106,7 @@ APPLY_SPECIFIC(conv_gi)(PyGpuArrayObject *kerns, PyGpuArrayObject *output,
     algo = choice.algo;
 #else
     size_t free;
-    int err2 = c->ops->property(c->ctx, NULL, NULL, GA_CTX_PROP_FREE_GMEM, &free);
+    int err2 = gpucontext_property(c->ctx, GA_CTX_PROP_FREE_GMEM, &free);
 
     if (err2 != GA_NO_ERROR) {
       PyErr_Format(PyExc_RuntimeError, "Error when trying to find the "
@@ -204,7 +204,7 @@ APPLY_SPECIFIC(conv_gi)(PyGpuArrayObject *kerns, PyGpuArrayObject *output,
   }
 
   if (worksize != 0) {
-    workspace = c->ops->buffer_alloc(c->ctx, worksize, NULL, 0, NULL);
+    workspace = gpudata_alloc(c->ctx, worksize, NULL, 0, NULL);
     if (workspace == NULL) {
       PyErr_SetString(PyExc_RuntimeError,
                       "Could not allocate working memory");
@@ -227,7 +227,7 @@ APPLY_SPECIFIC(conv_gi)(PyGpuArrayObject *kerns, PyGpuArrayObject *output,
     APPLY_SPECIFIC(input), PyGpuArray_DEV_DATA(*input));
 
   if (worksize != 0)
-    c->ops->buffer_release(workspace);
+    gpudata_release(workspace);
 
   cuda_record(kerns->ga.data, GPUARRAY_CUDA_WAIT_READ);
   cuda_record(output->ga.data, GPUARRAY_CUDA_WAIT_READ);

theano/gpuarray/dnn_gw.c

@@ -107,7 +107,7 @@ APPLY_SPECIFIC(conv_gw)(PyGpuArrayObject *input, PyGpuArrayObject *output,
     algo = choice.algo;
 #else
     size_t free;
-    int err2 = c->ops->property(c->ctx, NULL, NULL, GA_CTX_PROP_FREE_GMEM, &free);
+    int err2 = gpucontext_property(c->ctx, GA_CTX_PROP_FREE_GMEM, &free);
 
     if (err2 != GA_NO_ERROR) {
       PyErr_Format(PyExc_RuntimeError, "Error when trying to find the "
@@ -192,7 +192,7 @@ APPLY_SPECIFIC(conv_gw)(PyGpuArrayObject *input, PyGpuArrayObject *output,
   }
 
   if (worksize != 0) {
-    workspace = c->ops->buffer_alloc(c->ctx, worksize, NULL, 0, NULL);
+    workspace = gpudata_alloc(c->ctx, worksize, NULL, 0, NULL);
     if (workspace == NULL) {
       PyErr_SetString(PyExc_RuntimeError, "Could not allocate working memory");
       cuda_exit(c->ctx);
@@ -214,7 +214,7 @@ APPLY_SPECIFIC(conv_gw)(PyGpuArrayObject *input, PyGpuArrayObject *output,
     APPLY_SPECIFIC(kerns), PyGpuArray_DEV_DATA(*kerns));
 
   if (worksize != 0)
-    c->ops->buffer_release(workspace);
+    gpudata_release(workspace);
 
   cuda_record(input->ga.data, GPUARRAY_CUDA_WAIT_READ);
   cuda_record(output->ga.data, GPUARRAY_CUDA_WAIT_READ);

theano/gpuarray/extra_ops.py

@@ -249,17 +249,17 @@ class GpuCumsum(GpuKernelBase, Op):
                 size_t max_grid_size1;
                 size_t max_grid_size2;
                 int err;
-                err = %(ctx)s->ops->property(%(ctx)s->ctx, NULL, NULL, GA_CTX_PROP_MAXLSIZE0, &max_threads_dim0);
+                err = gpucontext_property(%(ctx)s->ctx, GA_CTX_PROP_MAXLSIZE0, &max_threads_dim0);
                 if (err != GA_NO_ERROR){
                     PyErr_SetString(PyExc_RuntimeError, "Could not fetch max_threads_dims0");
                     %(fail)s;
                 }
-                err = %(ctx)s->ops->property(%(ctx)s->ctx, NULL, NULL, GA_CTX_PROP_MAXGSIZE1, &max_grid_size1);
+                err = gpucontext_property(%(ctx)s->ctx, GA_CTX_PROP_MAXGSIZE1, &max_grid_size1);
                 if (err != GA_NO_ERROR){
                     PyErr_SetString(PyExc_RuntimeError, "Could not fetch max_grid_size1");
                     %(fail)s;
                 }
-                err = %(ctx)s->ops->property(%(ctx)s->ctx, NULL, NULL, GA_CTX_PROP_MAXGSIZE2, &max_grid_size2);
+                err = gpucontext_property(%(ctx)s->ctx, GA_CTX_PROP_MAXGSIZE2, &max_grid_size2);
                 if (err != GA_NO_ERROR){
                     PyErr_SetString(PyExc_RuntimeError, "Could not fetch max_grid_size2");
                     %(fail)s;

theano/gpuarray/gemm16.c

@@ -117,7 +117,7 @@ int gemm16(PyGpuArrayObject *C, float alpha,
         if (48 < n128 && n128 <= 64) {
           n64 = n / 64;
           if (nprocs == 0)
-            if (A->ga.ops->property(A->context->ctx, NULL, NULL,
+            if (gpucontext_property(A->context->ctx,
                                     GA_CTX_PROP_NUMPROCS, &nprocs)) {
               nprocs = 0;
               res = 1;

theano/gpuarray/subtensor.py

@@ -757,8 +757,8 @@ __device__ ga_half atomicAdd(ga_half *addr, ga_half val) {
             int err, kerr = 0;
 
             if (threads_per_block[0] > 0 && n_blocks[0] > 0) {
-              err = py_self->ga.ops->property(NULL, py_self->ga.data, NULL,
-                                              GA_CTX_PROP_ERRBUF, &errbuf);
+              err = gpudata_property(py_self->ga.data,
+                                     GA_CTX_PROP_ERRBUF, &errbuf);
               if (err != GA_NO_ERROR) {
                 PyErr_SetString(PyExc_RuntimeError, "Can't fetch error buffer");
                 return 1;
@@ -793,7 +793,7 @@ __device__ ga_half atomicAdd(ga_half *addr, ga_half val) {
                              GpuKernel_error(&%(k_var)s, err));
                 return 1;
               }
-              err = py_self->ga.ops->buffer_read(&kerr, errbuf, 0, sizeof(int));
+              err = gpudata_read(&kerr, errbuf, 0, sizeof(int));
               if (err != GA_NO_ERROR) {
                 PyErr_SetString(PyExc_RuntimeError, "Can't read error buffer");
                 return 1;
@@ -801,7 +801,7 @@ __device__ ga_half atomicAdd(ga_half *addr, ga_half val) {
               if (kerr != 0) {
                 PyErr_SetString(PyExc_IndexError, "Index out of bounds");
                 kerr = 0;
-                py_self->ga.ops->buffer_write(errbuf, 0, &kerr, sizeof(int));
+                gpudata_write(errbuf, 0, &kerr, sizeof(int));
                 return 1;
               }
             }