Can you let us know which compiler (and which version) you are using?

On Thu, Sep 22, 2016, D. Kit wrote:
> Hello,
> 
> I'm attempting to use conv3D2D in Theano, but when I get to the function 
> compilation stage it fails with the following errors:
> 
> mod.cpp:684:1: error: expected unqualified-id before '{' token
> 
> mod.cpp:1345:1: error: expected '}' at end of input
> mod.cpp: In destructor '{anonymous}::__struct_compiled_op_d7db26d1cc9f
> 
> mod.cpp:573:21: error: 'cleanup' was not declared in this scope
> mod.cpp: At global scope:
> mod.cpp:678:1: error: expected unqualified-id at end of input
>  #undef ITEMSIZE_INPUT_2
>  ^
> mod.cpp:678:1: error: expected '}' at end of input
> 
> I'm attaching the mod.cpp file that was generated by cc.py, op.py, dnn.py, 
> etc., The initialization code from dnn_base.c starts at line 684, but there 
> is no function name defined, so the code sits outside any useful code 
> block. The "#section cleanup_code_struct" code section also appears to have 
> been inserted into the wrong spot.
> 
> Can someone tell me how to fix this?
> 
> --Dmitry
> 
> -- 
> 
> --- 
> You received this message because you are subscribed to the Google Groups 
> "theano-users" group.
> To unsubscribe from this group and stop receiving emails from it, send an 
> email to theano-users+unsubscr...@googlegroups.com.
> For more options, visit https://groups.google.com/d/optout.

> #include <Python.h>
> #include <iostream>
> #include "theano_mod_helper.h"
> #include <gpuarray/array.h>
> #include <gpuarray/kernel.h>
> #include <gpuarray/error.h>
> #include <gpuarray/buffer.h>
> #include <gpuarray/buffer_blas.h>
> #include <numpy/arrayobject.h>
> #include <gpuarray_api.h>
> #include <math.h>
> #include <numpy/arrayscalars.h>
> #include "cudnn.h"
> #include "cudnn_helper.h"
> #include "gpuarray_helper.h"
> #include "gpuarray/types.h"
> #include "gpuarray/array.h"
> #include "gpuarray/util.h"
> #include "gpuarray/ext_cuda.h"
> #include "gpuarray_api.h"
> #include "numpy_compat.h"
> //////////////////////
> ////  Support Code
> //////////////////////
> 
> void _capsule_destructor(PyObject *o) {
>     void *d = PyCapsule_GetContext(o);
>     void *p = PyCapsule_GetPointer(o, NULL);
>     void (*f)(void *) = (void (*)(void *))d;
>     if (f != NULL) f(p);
> }
> 
> 
> 
> 
> static int
> c_set_tensorNd(PyGpuArrayObject *var, cudnnTensorDescriptor_t desc) {
>   cudnnDataType_t dt;
>   size_t ds;
>   switch (var->ga.typecode) {
>   case GA_FLOAT:
>     dt = CUDNN_DATA_FLOAT;
>     break;
>   case GA_DOUBLE:
>     dt = CUDNN_DATA_DOUBLE;
>     break;
>   case GA_HALF:
>     dt = CUDNN_DATA_HALF;
>     break;
>   default:
>     PyErr_SetString(PyExc_TypeError, "Non-float datatype in c_set_tensorNd");
>     return -1;
>   }
>   ds = gpuarray_get_elsize(var->ga.typecode);
> 
>   int strs[8], dims[8], default_stride = 1;
>   unsigned int nd = PyGpuArray_NDIM(var);
> 
>   if (nd > 8) {
>     PyErr_SetString(PyExc_TypeError, "Tensor of more than 8d");
>     return -1;
>   }
> 
>   for (unsigned int _i = nd; _i > 0; _i--) {
>     unsigned int i = _i - 1;
>     strs[i] = PyGpuArray_STRIDE(var, i) ?
>       PyGpuArray_STRIDE(var, i)/ds : default_stride;
>     default_stride *= PyGpuArray_DIM(var, i);
>     dims[i] = PyGpuArray_DIM(var, i);
>   }
> 
>   /* Tensors can't be smaller than 3d for cudnn so we pad the
>    * descriptor if they are */
>   for (unsigned int i = nd; i < 3; i++) {
>     strs[i] = 1;
>     dims[i] = 1;
>   }
> 
>   cudnnStatus_t err = cudnnSetTensorNdDescriptor(desc, dt, nd < 3 ? 3 : nd,
>                                                  dims, strs);
>   if (err != CUDNN_STATUS_SUCCESS) {
>     PyErr_Format(PyExc_RuntimeError,
>                "Could not set tensorNd descriptor: %s",
>                cudnnGetErrorString(err));
>     return -1;
>   }
>   return 0;
> }
> 
> static int c_make_tensorNd(PyGpuArrayObject *var, cudnnTensorDescriptor_t 
> *desc) {
>   cudnnStatus_t err;
>   err = cudnnCreateTensorDescriptor(desc);
>   if (err != CUDNN_STATUS_SUCCESS) {
>     PyErr_Format(PyExc_RuntimeError,
>                  "Could not create tensor descriptor: %s",
>                  cudnnGetErrorString(err));
>     return -1;
>   }
>   if (c_set_tensorNd(var, *desc) != 0) {
>     cudnnDestroyTensorDescriptor(*desc);
>     return -1;
>   }
>   return 0;
> }
> 
> static int
> c_set_filter(PyGpuArrayObject *var, cudnnFilterDescriptor_t desc) {
>   cudnnDataType_t dt;
>   cudnnStatus_t err;
> 
>   if (!GpuArray_IS_C_CONTIGUOUS(&var->ga)) {
>     PyErr_SetString(PyExc_ValueError,
>                   "Only contiguous filters (kernels) are supported.");
>     return -1;
>   }
>   switch (var->ga.typecode) {
>   case GA_FLOAT:
>     dt = CUDNN_DATA_FLOAT;
>     break;
>   case GA_DOUBLE:
>     dt = CUDNN_DATA_DOUBLE;
>     break;
>   case GA_HALF:
>     dt = CUDNN_DATA_HALF;
>     break;
>   default:
>     PyErr_SetString(PyExc_TypeError, "Non-float datatype in c_set_filter");
>     return -1;
>   }
> 
>   int dims[8];
>   unsigned int nd = PyGpuArray_NDIM(var);
> 
>   if (nd > 8) {
>     PyErr_SetString(PyExc_TypeError, "Tensor of more than 8d");
>     return -1;
>   }
> 
>   for (unsigned int _i = nd; _i > 0; _i--) {
>     unsigned int i = _i - 1;
>     dims[i] = PyGpuArray_DIM(var, i);
>   }
> 
>   /* Filters can't be less than 3d so we pad */
>   for (unsigned int i = nd; i < 3; i++)
>     dims[i] = 1;
> 
>   if (nd < 3)
>     nd = 3;
> 
> #if CUDNN_VERSION >= 5000
>     err = cudnnSetFilterNdDescriptor(desc, dt, CUDNN_TENSOR_NCHW, nd, dims);
> #else
>     err = cudnnSetFilterNdDescriptor(desc, dt, nd, dims);
> #endif
> 
>   if (err != CUDNN_STATUS_SUCCESS) {
>     PyErr_Format(PyExc_RuntimeError,
>                "Could not set filter descriptor: %s.",
>                cudnnGetErrorString(err));
>     return -1;
>   }
>   return 0;
> }
> 
> static int c_make_filter(PyGpuArrayObject *var, cudnnFilterDescriptor_t 
> *desc) {
>   cudnnStatus_t err;
>   err = cudnnCreateFilterDescriptor(desc);
>   if (err != CUDNN_STATUS_SUCCESS) {
>     PyErr_Format(PyExc_RuntimeError,
>                  "Could not create tensor descriptor: %s",
>                  cudnnGetErrorString(err));
>     return -1;
>   }
>   if (c_set_filter(var, *desc) != 0) {
>     cudnnDestroyFilterDescriptor(*desc);
>     return -1;
>   }
>   return 0;
> }
> 
> 
> 
>     namespace {
>     struct __struct_compiled_op_d7db26d1cc9fb59ee56289f4dda72577 {
>         PyObject* __ERROR;
> 
>         PyObject* storage_V3;
> PyObject* storage_V5;
> PyObject* storage_V7;
> PyObject* storage_V9;
> PyObject* storage_V11;
> PyObject* storage_V13;
> PyObject* storage_V1;
> PyObject* storage_V15;
>         
>     PyObject* py_V15;
>     PyGpuContextObject *V15;
> 
> #define DTYPE_INPUT_0 npy_float32
> 
> #define TYPENUM_INPUT_0 11
> 
> #define ITEMSIZE_INPUT_0 4
> 
> #define DTYPE_INPUT_1 npy_float32
> 
> #define TYPENUM_INPUT_1 11
> 
> #define ITEMSIZE_INPUT_1 4
> 
> #define DTYPE_INPUT_2 npy_float32
> 
> #define TYPENUM_INPUT_2 11
> 
> #define ITEMSIZE_INPUT_2 4
> 
> #define DTYPE_INPUT_4 npy_float32
> 
> #define TYPENUM_INPUT_4 11
> 
> #define ITEMSIZE_INPUT_4 4
> 
> #define DTYPE_INPUT_5 npy_float32
> 
> #define TYPENUM_INPUT_5 11
> 
> #define ITEMSIZE_INPUT_5 4
> 
> #define DTYPE_OUTPUT_0 npy_float32
> 
> #define TYPENUM_OUTPUT_0 11
> 
> #define ITEMSIZE_OUTPUT_0 4
> 
> #define APPLY_SPECIFIC(str) str##_node_d7db26d1cc9fb59ee56289f4dda72577_0
> 
> #define CONV_INPLACE 1
> 
> #define CONV_ALGO CUDNN_CONVOLUTION_FWD_ALGO_IMPLICIT_PRECOMP_GEMM
> 
> 
> 
> PyGpuContextObject *ctx;
> cudnnHandle_t APPLY_SPECIFIC(_handle);
> 
> 
> cudnnTensorDescriptor_t APPLY_SPECIFIC(input);
> cudnnTensorDescriptor_t APPLY_SPECIFIC(output);
> cudnnFilterDescriptor_t APPLY_SPECIFIC(kerns);
> 
> 
> 
> #ifdef CHOOSE_ALGO
> int reuse_algo;
> cudnnConvolutionFwdAlgo_t prev_algo;
> #ifndef CHOOSE_ONCE
> size_t prev_img_dims[5];
> size_t prev_kern_dims[5];
> #endif
> #endif
> 
> int
> APPLY_SPECIFIC(conv_fwd)(PyGpuArrayObject *input, PyGpuArrayObject *kerns,
>                          PyGpuArrayObject *om,
>                          cudnnConvolutionDescriptor_t desc,
>                          double alpha, double beta,
>                          PyGpuArrayObject **output,
>                          PyGpuContextObject *c) {
>   cudnnStatus_t err = CUDNN_STATUS_SUCCESS;
>   float af = alpha, bf = beta;
>   void *alpha_p;
>   void *beta_p;
> 
>   if (PyGpuArray_DIMS(input)[1] != PyGpuArray_DIMS(kerns)[1]) {
>     PyErr_SetString(PyExc_ValueError,
>                   "images and kernel must have the same stack size");
>     return 1;
>   }
> 
>   if (c_set_tensorNd(input, APPLY_SPECIFIC(input)) == -1)
>     return 1;
>   if (c_set_filter(kerns, APPLY_SPECIFIC(kerns)) == -1)
>     return 1;
> 
>   switch (input->ga.typecode) {
>   case GA_DOUBLE:
>     alpha_p = (void *)&alpha;
>     beta_p = (void *)&beta;
>     break;
>   case GA_FLOAT:
>   case GA_HALF:
>     alpha_p = (void *)&af;
>     beta_p = (void *)&bf;
>     break;
>   default:
>     PyErr_SetString(PyExc_TypeError, "Unsupported type in convolution");
>     return 1;
>   }
> 
> #ifdef CONV_INPLACE
>   Py_XDECREF(*output);
>   *output = om;
>   Py_INCREF(*output);
> #else
>   if (theano_prep_output(output, PyGpuArray_NDIM(om), PyGpuArray_DIMS(om),
>                          om->ga.typecode, GA_C_ORDER, c) != 0)
>     return 1;
>   if (beta != 0.0 && pygpu_move(*output, om))
>     return 1;
> #endif
> 
>   if (c_set_tensorNd(*output, APPLY_SPECIFIC(output)) == -1)
>     return 1;
> 
>   cudnnConvolutionFwdAlgo_t algo = CONV_ALGO;
> 
>   cuda_enter(c->ctx);
> #ifdef CHOOSE_ALGO
> #ifndef CHOOSE_ONCE
>   reuse_algo = 1;
>   for (unsigned int i = 0; i < PyGpuArray_NDIM(input); i++) {
>     reuse_algo = (reuse_algo &&
>                   PyGpuArray_DIM(input, i) == prev_img_dims[i]);
>     reuse_algo = (reuse_algo &&
>                   PyGpuArray_DIM(kerns, i) == prev_kern_dims[i]);
>   }
> #endif
> 
>   if (!reuse_algo) {
> #ifdef CHOOSE_TIME
>     int count;
>     cudnnConvolutionFwdAlgoPerf_t choice;
>     err = cudnnFindConvolutionForwardAlgorithm(
>       APPLY_SPECIFIC(_handle), APPLY_SPECIFIC(input), APPLY_SPECIFIC(kerns),
>       desc, APPLY_SPECIFIC(output), 1, &count, &choice);
> 
>     if (err != CUDNN_STATUS_SUCCESS) {
>       PyErr_Format(PyExc_RuntimeError,
>                    "error selecting convolution algo: %s",
>                    cudnnGetErrorString(err));
>       cuda_exit(c->ctx);
>       return 1;
>     }
>     algo = choice.algo;
> #else
>     size_t 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 "
>                    "memory information on the GPU");
>       cuda_exit(c->ctx);
>       return 1;
>     }
> 
>     err = cudnnGetConvolutionForwardAlgorithm(
>       APPLY_SPECIFIC(_handle), APPLY_SPECIFIC(input), APPLY_SPECIFIC(kerns),
>       desc, APPLY_SPECIFIC(output),
>       CUDNN_CONVOLUTION_FWD_SPECIFY_WORKSPACE_LIMIT, free, &algo);
>     if (err != CUDNN_STATUS_SUCCESS) {
>       PyErr_Format(PyExc_RuntimeError,
>                    "error selecting convolution algo: %s",
>                    cudnnGetErrorString(err));
>       cuda_exit(c->ctx);
>       return 1;
>     }
> #endif
>     prev_algo = algo;
>   } else {
>     algo = prev_algo;
>   }
> 
> #ifdef CHOOSE_ONCE
>   reuse_algo = 1;
> #else
>   for (unsigned int i = 0; i < PyGpuArray_NDIM(input); i++) {
>     prev_img_dims[i] = PyGpuArray_DIM(input, i);
>     prev_kern_dims[i] = PyGpuArray_DIM(kerns, i);
>   }
> #endif
> 
> #endif
> 
>   /* These two algos are not supported for 3d conv */
>   if (PyGpuArray_NDIM(input) == 5 &&
>       (algo == CUDNN_CONVOLUTION_FWD_ALGO_IMPLICIT_PRECOMP_GEMM ||
>        algo == CUDNN_CONVOLUTION_FWD_ALGO_GEMM))
>     algo = CUDNN_CONVOLUTION_FWD_ALGO_IMPLICIT_GEMM;
> 
>   // The FFT implementation does not support strides, 1x1 filters or inputs
>   // with a spatial dimension larger than 1024. The tiled-FFT implementation
>   // does not support strides.
>   // If the chosen implementation is FFT or tiled-FFT, validate that it can
>   // be used on the current data and default to a safe implementation if it
>   // can't.
>   // The following code is 2d-specific but it is fine as FFT and tiled-FFT are
>   // defined only for 2d filters
>   if ((algo == CUDNN_CONVOLUTION_FWD_ALGO_FFT ||
>        algo == CUDNN_CONVOLUTION_FWD_ALGO_FFT_TILING) && 
> PyGpuArray_NDIM(input) == 4) {
> 
>     // Extract the properties of the convolution descriptor
>     int nd;
>     int pad[2];
>     int stride[2];
>     int upscale[2];
>     cudnnConvolutionMode_t mode;
>     cudnnDataType_t data_type;
>     err = cudnnGetConvolutionNdDescriptor(desc, 2, &nd, pad, stride,
>                                              upscale, &mode, &data_type);
>     if (err != CUDNN_STATUS_SUCCESS) {
>       PyErr_Format(PyExc_RuntimeError,
>                    "error getting convolution properties: %s",
>                    cudnnGetErrorString(err));
>       cuda_exit(c->ctx);
>       return 1;
>     }
> 
>     if (algo == CUDNN_CONVOLUTION_FWD_ALGO_FFT)
>     {
>       if (stride[0] != 1 || stride[1] != 1 ||
>           PyGpuArray_DIM(input, 2) > 1024 || PyGpuArray_DIM(input, 3) > 1024 
> ||
>           (PyGpuArray_DIM(kerns, 2) == 1 && PyGpuArray_DIM(kerns, 3) == 1))
>       {
>         algo = CUDNN_CONVOLUTION_FWD_ALGO_IMPLICIT_GEMM;
>       }
>     }
>     else
>     {
>       // algo == CUDNN_CONVOLUTION_FWD_ALGO_FFT_TILING
>       if (stride[0] != 1 || stride[1] != 1)
>       {
>         algo = CUDNN_CONVOLUTION_FWD_ALGO_IMPLICIT_GEMM;
>       }
>     }
>   }
> 
>   {
>     size_t worksize;
>     gpudata *workspace;
>     err = cudnnGetConvolutionForwardWorkspaceSize(APPLY_SPECIFIC(_handle),
>                                                   APPLY_SPECIFIC(input),
>                                                   APPLY_SPECIFIC(kerns),
>                                                   desc,
>                                                   APPLY_SPECIFIC(output),
>                                                   algo,
>                                                   &worksize);
> 
>     if (err == CUDNN_STATUS_NOT_SUPPORTED) {
>       // Fallback to none algo if not supported
>       // TODO: Print a warning
>       algo = CUDNN_CONVOLUTION_FWD_ALGO_IMPLICIT_GEMM;
> 
>       err = cudnnGetConvolutionForwardWorkspaceSize(APPLY_SPECIFIC(_handle),
>                                                     APPLY_SPECIFIC(input),
>                                                     APPLY_SPECIFIC(kerns),
>                                                     desc,
>                                                     APPLY_SPECIFIC(output),
>                                                     algo,
>                                                     &worksize);
>     }
> 
>     if (err != CUDNN_STATUS_SUCCESS) {
>       PyErr_Format(PyExc_RuntimeError,
>                    "error getting worksize: %s",
>                    cudnnGetErrorString(err));
>       cuda_exit(c->ctx);
>       return 1;
>     }
> 
>     /*
>      * This is less than ideal since we need to free it after (which
>      * introduces a synchronization point. But we don't have a module
>      * to place a nice get_work_mem() function in.
>      */
>     if (worksize != 0) {
>       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);
>         return 1;
>       }
>     }
> 
>     cuda_wait(input->ga.data, GPUARRAY_CUDA_WAIT_READ);
>     cuda_wait(kerns->ga.data, GPUARRAY_CUDA_WAIT_READ);
>     cuda_wait((*output)->ga.data, GPUARRAY_CUDA_WAIT_WRITE);
> 
>     err = cudnnConvolutionForward(
>       APPLY_SPECIFIC(_handle),
>       alpha_p,
>       APPLY_SPECIFIC(input), PyGpuArray_DEV_DATA(input),
>       APPLY_SPECIFIC(kerns), PyGpuArray_DEV_DATA(kerns),
>       desc, algo,
>       worksize == 0 ? NULL : *(void **)workspace, worksize,
>       beta_p,
>       APPLY_SPECIFIC(output), PyGpuArray_DEV_DATA(*output));
> 
>     if (worksize != 0)
>       gpudata_release(workspace);
> 
>     cuda_record(input->ga.data, GPUARRAY_CUDA_WAIT_READ);
>     cuda_record(kerns->ga.data, GPUARRAY_CUDA_WAIT_READ);
>     cuda_record((*output)->ga.data, GPUARRAY_CUDA_WAIT_WRITE);
>   }
>   cuda_exit(c->ctx);
> 
>   if (err != CUDNN_STATUS_SUCCESS) {
>     PyErr_Format(PyExc_RuntimeError, "error doing operation: %s",
>                cudnnGetErrorString(err));
>     return 1;
>   }
>   return 0;
> }
> 
> 
> #undef DTYPE_INPUT_0
> 
> #undef TYPENUM_INPUT_0
> 
> #undef ITEMSIZE_INPUT_0
> 
> #undef DTYPE_INPUT_1
> 
> #undef TYPENUM_INPUT_1
> 
> #undef ITEMSIZE_INPUT_1
> 
> #undef DTYPE_INPUT_2
> 
> #undef TYPENUM_INPUT_2
> 
> #undef ITEMSIZE_INPUT_2
> 
> #undef DTYPE_INPUT_4
> 
> #undef TYPENUM_INPUT_4
> 
> #undef ITEMSIZE_INPUT_4
> 
> #undef DTYPE_INPUT_5
> 
> #undef TYPENUM_INPUT_5
> 
> #undef ITEMSIZE_INPUT_5
> 
> #undef DTYPE_OUTPUT_0
> 
> #undef TYPENUM_OUTPUT_0
> 
> #undef ITEMSIZE_OUTPUT_0
> 
> #undef APPLY_SPECIFIC
> 
> #undef CONV_INPLACE
> 
> #undef CONV_ALGO
> 
>         __struct_compiled_op_d7db26d1cc9fb59ee56289f4dda72577() {
>             // This is only somewhat safe because we:
>             //  1) Are not a virtual class
>             //  2) Do not use any virtual classes in the members
>             //  3) Deal with mostly POD and pointers
> 
>             // If this changes, we would have to revise this, but for
>             // now I am tired of chasing segfaults because
>             // initialization code had an error and some pointer has
>             // a junk value.
>             memset(this, 0, sizeof(*this));
>         }
>         ~__struct_compiled_op_d7db26d1cc9fb59ee56289f4dda72577(void) {
>             cleanup();
>         }
> 
>         int init(PyObject* __ERROR, PyObject* storage_V3, PyObject* 
> storage_V5, PyObject* storage_V7, PyObject* storage_V9, PyObject* 
> storage_V11, PyObject* storage_V13, PyObject* storage_V1, PyObject* 
> storage_V15) {
>             Py_XINCREF(storage_V3);
> Py_XINCREF(storage_V5);
> Py_XINCREF(storage_V7);
> Py_XINCREF(storage_V9);
> Py_XINCREF(storage_V11);
> Py_XINCREF(storage_V13);
> Py_XINCREF(storage_V1);
> Py_XINCREF(storage_V15);
>             this->storage_V3 = storage_V3;
> this->storage_V5 = storage_V5;
> this->storage_V7 = storage_V7;
> this->storage_V9 = storage_V9;
> this->storage_V11 = storage_V11;
> this->storage_V13 = storage_V13;
> this->storage_V1 = storage_V1;
> this->storage_V15 = storage_V15;
>             
> 
> 
> 
> 
> 
> 
> 
> 
>     py_V15 = PyList_GET_ITEM(storage_V15, 0);
>     {Py_XINCREF(py_V15);}
>     
> if (!PyObject_TypeCheck(py_V15, &PyGpuContextType)) {
>   PyErr_SetString(PyExc_TypeError, "expected a GpuContext");
>   {
>         if (!PyErr_Occurred()) {
>             PyErr_SetString(PyExc_RuntimeError,
>                 "Unexpected error in an Op's C code. "
>                 "No Python exception was set.");
>             }
>         return 15;
> }
> }
> 
> V15 = (PyGpuContextObject *)py_V15;
> Py_INCREF(V15);
> 
> 
> 
> #define DTYPE_INPUT_0 npy_float32
> 
> #define TYPENUM_INPUT_0 11
> 
> #define ITEMSIZE_INPUT_0 4
> 
> #define DTYPE_INPUT_1 npy_float32
> 
> #define TYPENUM_INPUT_1 11
> 
> #define ITEMSIZE_INPUT_1 4
> 
> #define DTYPE_INPUT_2 npy_float32
> 
> #define TYPENUM_INPUT_2 11
> 
> #define ITEMSIZE_INPUT_2 4
> 
> #define DTYPE_INPUT_4 npy_float32
> 
> #define TYPENUM_INPUT_4 11
> 
> #define ITEMSIZE_INPUT_4 4
> 
> #define DTYPE_INPUT_5 npy_float32
> 
> #define TYPENUM_INPUT_5 11
> 
> #define ITEMSIZE_INPUT_5 4
> 
> #define DTYPE_OUTPUT_0 npy_float32
> 
> #define TYPENUM_OUTPUT_0 11
> 
> #define ITEMSIZE_OUTPUT_0 4
> 
> #define APPLY_SPECIFIC(str) str##_node_d7db26d1cc9fb59ee56289f4dda72577_0
> 
> #define CONV_INPLACE 1
> 
> #define CONV_ALGO CUDNN_CONVOLUTION_FWD_ALGO_IMPLICIT_PRECOMP_GEMM
> 
> #define FAIL { \
> 
>         if (!PyErr_Occurred()) { \
> 
>             PyErr_SetString(PyExc_RuntimeError, \
> 
>                 "Unexpected error in an Op's C code. " \
> 
>                 "No Python exception was set."); \
> 
>             } \
> 
>         return 17; \
> 
> }
> 
> #define PARAMS V15
> 
> 
> 
> {
>   // We need to keep a reference here to have it available in the destructor.
>   ctx = PARAMS;
>   Py_INCREF(ctx);
> 
>   cuda_enter(PARAMS->ctx);
>   cudnnStatus_t err;
>   APPLY_SPECIFIC(_handle) = NULL;
>   if ((err = cudnnCreate(&APPLY_SPECIFIC(_handle))) != CUDNN_STATUS_SUCCESS) {
>     PyErr_Format(PyExc_RuntimeError, "could not create cuDNN handle: %s",
>                  cudnnGetErrorString(err));
>     cuda_exit(PARAMS->ctx);
>     FAIL;
>   }
>   if ((err = cudnnSetStream(APPLY_SPECIFIC(_handle),
>                             cuda_get_stream(PARAMS->ctx))) != 
> CUDNN_STATUS_SUCCESS) {
>     PyErr_Format(PyExc_RuntimeError, "Could not set cudnn stream: %s",
>                  cudnnGetErrorString(err));
>     cuda_exit(PARAMS->ctx);
>     FAIL;
>   }
>   cuda_exit(PARAMS->ctx);
> }
> 
> 
> 
> cudnnStatus_t APPLY_SPECIFIC(err);
> APPLY_SPECIFIC(input) = NULL;
> APPLY_SPECIFIC(output) = NULL;
> APPLY_SPECIFIC(kerns) = NULL;
> if ((APPLY_SPECIFIC(err) = 
> cudnnCreateTensorDescriptor(&APPLY_SPECIFIC(input))) != CUDNN_STATUS_SUCCESS) 
> {
>   PyErr_Format(PyExc_MemoryError, "could not allocate tensor descriptor "
>              "(inp): %s", cudnnGetErrorString(APPLY_SPECIFIC(err)));
>   FAIL;
> }
> if ((APPLY_SPECIFIC(err) = 
> cudnnCreateTensorDescriptor(&APPLY_SPECIFIC(output))) != 
> CUDNN_STATUS_SUCCESS) {
>   PyErr_Format(PyExc_MemoryError, "could not allocate tensor descriptor "
>                "(out): %s", cudnnGetErrorString(APPLY_SPECIFIC(err)));
>   FAIL;
> }
> if ((APPLY_SPECIFIC(err) = 
> cudnnCreateFilterDescriptor(&APPLY_SPECIFIC(kerns))) != CUDNN_STATUS_SUCCESS) 
> {
>   PyErr_Format(PyExc_MemoryError, "could not allocate filter descriptor: %s", 
>              cudnnGetErrorString(APPLY_SPECIFIC(err)));
>   FAIL;
> }
> 
> 
> 
> #ifdef CHOOSE_ALGO
> reuse_algo = 0;
> prev_algo = CONV_ALGO;
> #ifndef CHOOSE_ONCE
> memset(prev_img_dims, 0, sizeof(prev_img_dims));
> memset(prev_kern_dims, 0, sizeof(prev_kern_dims));
> #endif
> #endif
> 
> 
> 
> #undef FAIL
> 
> #undef PARAMS
> 
> #undef DTYPE_INPUT_0
> 
> #undef TYPENUM_INPUT_0
> 
> #undef ITEMSIZE_INPUT_0
> 
> #undef DTYPE_INPUT_1
> 
> #undef TYPENUM_INPUT_1
> 
> #undef ITEMSIZE_INPUT_1
> 
> #undef DTYPE_INPUT_2
> 
> #undef TYPENUM_INPUT_2
> 
> #undef ITEMSIZE_INPUT_2
> 
> #undef DTYPE_INPUT_4
> 
> #undef TYPENUM_INPUT_4
> 
> #undef ITEMSIZE_INPUT_4
> 
> #undef DTYPE_INPUT_5
> 
> #undef TYPENUM_INPUT_5
> 
> #undef ITEMSIZE_INPUT_5
> 
> #undef DTYPE_OUTPUT_0
> 
> #undef TYPENUM_OUTPUT_0
> 
> #undef ITEMSIZE_OUTPUT_0
> 
> #undef APPLY_SPECIFIC
> 
> #undef CONV_INPLACE
> 
> #undef CONV_ALGO
>             this->__ERROR = __ERROR;
>             return 0;
>         }
>         void cleanup(void) {
>             __label_1:
> 
> double __DUMMY_1;
> __label_3:
> 
> double __DUMMY_3;
> __label_5:
> 
> double __DUMMY_5;
> __label_7:
> 
> double __DUMMY_7;
> __label_9:
> 
> double __DUMMY_9;
> __label_11:
> 
> double __DUMMY_11;
> __label_13:
> 
> double __DUMMY_13;
> __label_15:
> Py_XDECREF(V15); V15 = NULL;
>     {Py_XDECREF(py_V15);}
>     
> double __DUMMY_15;
> __label_18:
> 
> 
> #define DTYPE_INPUT_0 npy_float32
> 
> #define TYPENUM_INPUT_0 11
> 
> #define ITEMSIZE_INPUT_0 4
> 
> #define DTYPE_INPUT_1 npy_float32
> 
> #define TYPENUM_INPUT_1 11
> 
> #define ITEMSIZE_INPUT_1 4
> 
> #define DTYPE_INPUT_2 npy_float32
> 
> #define TYPENUM_INPUT_2 11
> 
> #define ITEMSIZE_INPUT_2 4
> 
> #define DTYPE_INPUT_4 npy_float32
> 
> #define TYPENUM_INPUT_4 11
> 
> #define ITEMSIZE_INPUT_4 4
> 
> #define DTYPE_INPUT_5 npy_float32
> 
> #define TYPENUM_INPUT_5 11
> 
> #define ITEMSIZE_INPUT_5 4
> 
> #define DTYPE_OUTPUT_0 npy_float32
> 
> #define TYPENUM_OUTPUT_0 11
> 
> #define ITEMSIZE_OUTPUT_0 4
> 
> #define APPLY_SPECIFIC(str) str##_node_d7db26d1cc9fb59ee56289f4dda72577_0
> 
> #define CONV_INPLACE 1
> 
> #define CONV_ALGO CUDNN_CONVOLUTION_FWD_ALGO_IMPLICIT_PRECOMP_GEMM
> 
> 
> 
> cuda_enter(ctx->ctx);
> cudnnDestroy(APPLY_SPECIFIC(_handle));
> cuda_exit(ctx->ctx);
> Py_DECREF((PyObject *)ctx);
> 
> 
> if (APPLY_SPECIFIC(input) != NULL)
>   cudnnDestroyTensorDescriptor(APPLY_SPECIFIC(input));
> if (APPLY_SPECIFIC(output) != NULL)
>   cudnnDestroyTensorDescriptor(APPLY_SPECIFIC(output));
> if (APPLY_SPECIFIC(kerns) != NULL)
>   cudnnDestroyFilterDescriptor(APPLY_SPECIFIC(kerns));
> 
> 
> #undef DTYPE_INPUT_0
> 
> #undef TYPENUM_INPUT_0
> 
> #undef ITEMSIZE_INPUT_0
> 
> #undef DTYPE_INPUT_1
> 
> #undef TYPENUM_INPUT_1
> 
> #undef ITEMSIZE_INPUT_1
> 
> #undef DTYPE_INPUT_2
> 
> #undef TYPENUM_INPUT_2
> 
> #undef ITEMSIZE_INPUT_2
> 
> #undef DTYPE_INPUT_4
> 
> #undef TYPENUM_INPUT_4
> 
> #undef ITEMSIZE_INPUT_4
> 
> #undef DTYPE_INPUT_5
> 
> #undef TYPENUM_INPUT_5
> 
> #undef ITEMSIZE_INPUT_5
> 
> #undef DTYPE_OUTPUT_0
> 
> #undef TYPENUM_OUTPUT_0
> 
> #undef ITEMSIZE_OUTPUT_0
> 
> #undef APPLY_SPECIFIC
> 
> #undef CONV_INPLACE
> 
> #undef CONV_ALGO
> double __DUMMY_18;
> 
>             Py_XDECREF(this->storage_V3);
> Py_XDECREF(this->storage_V5);
> Py_XDECREF(this->storage_V7);
> Py_XDECREF(this->storage_V9);
> Py_XDECREF(this->storage_V11);
> Py_XDECREF(this->storage_V13);
> Py_XDECREF(this->storage_V1);
> Py_XDECREF(this->storage_V15);
>         }
>         int run(void) {
>             int __failure = 0;
>             
>     PyObject* py_V1;
>     
>         PyGpuArrayObject *V1;
>         
>     PyObject* py_V3;
>     
>         PyGpuArrayObject *V3;
>         
>     PyObject* py_V5;
>     
>         PyGpuArrayObject *V5;
>         
>     PyObject* py_V7;
>     
>         PyGpuArrayObject *V7;
>         
>     PyObject* py_V9;
>     
>         cudnnConvolutionDescriptor_t V9;
>         
>     PyObject* py_V11;
>     
>                 typedef npy_float32 V11_dtype; // Deprecated use dtype_V11 
> instead.
>                 typedef npy_float32 dtype_V11;
>             
>         npy_float32 V11;
>         
>     PyObject* py_V13;
>     
>                 typedef npy_float32 V13_dtype; // Deprecated use dtype_V13 
> instead.
>                 typedef npy_float32 dtype_V13;
>             
>         npy_float32 V13;
>         
> {
> 
>     py_V1 = PyList_GET_ITEM(storage_V1, 0);
>     {Py_XINCREF(py_V1);}
>     
>         if (py_V1 == Py_None)
>         {
>             V1 = NULL;
>         }
>         else
>         {
>             
>         V1 = NULL;
>         if (py_V1 == Py_None) {
>             PyErr_SetString(PyExc_ValueError, "expected a GpuArray, not 
> None");
>             {
>         __failure = 2;
>         if (!PyErr_Occurred()) {
>             PyErr_SetString(PyExc_RuntimeError,
>                 "Unexpected error in an Op's C code. "
>                 "No Python exception was set.");
>             }
>         goto __label_2;}
>         }
>         /* First check if we are the base type exactly (the most common case),
>            then do the full subclass check if needed. */
>         if (py_V1->ob_type != &PyGpuArrayType &&
>             !PyObject_TypeCheck(py_V1, &PyGpuArrayType)) {
>             PyErr_SetString(PyExc_ValueError, "expected a GpuArray");
>             {
>         __failure = 2;
>         if (!PyErr_Occurred()) {
>             PyErr_SetString(PyExc_RuntimeError,
>                 "Unexpected error in an Op's C code. "
>                 "No Python exception was set.");
>             }
>         goto __label_2;}
>         }
>         V1 = (PyGpuArrayObject *)py_V1;
>         Py_INCREF(V1);
>         
>         }
>         
> {
> 
>     py_V3 = PyList_GET_ITEM(storage_V3, 0);
>     {Py_XINCREF(py_V3);}
>     
>         V3 = NULL;
>         if (py_V3 == Py_None) {
>             PyErr_SetString(PyExc_ValueError, "expected a GpuArray, not 
> None");
>             {
>         __failure = 4;
>         if (!PyErr_Occurred()) {
>             PyErr_SetString(PyExc_RuntimeError,
>                 "Unexpected error in an Op's C code. "
>                 "No Python exception was set.");
>             }
>         goto __label_4;}
>         }
>         /* First check if we are the base type exactly (the most common case),
>            then do the full subclass check if needed. */
>         if (py_V3->ob_type != &PyGpuArrayType &&
>             !PyObject_TypeCheck(py_V3, &PyGpuArrayType)) {
>             PyErr_SetString(PyExc_ValueError, "expected a GpuArray");
>             {
>         __failure = 4;
>         if (!PyErr_Occurred()) {
>             PyErr_SetString(PyExc_RuntimeError,
>                 "Unexpected error in an Op's C code. "
>                 "No Python exception was set.");
>             }
>         goto __label_4;}
>         }
>         V3 = (PyGpuArrayObject *)py_V3;
>         Py_INCREF(V3);
>         
> {
> 
>     py_V5 = PyList_GET_ITEM(storage_V5, 0);
>     {Py_XINCREF(py_V5);}
>     
>         V5 = NULL;
>         if (py_V5 == Py_None) {
>             PyErr_SetString(PyExc_ValueError, "expected a GpuArray, not 
> None");
>             {
>         __failure = 6;
>         if (!PyErr_Occurred()) {
>             PyErr_SetString(PyExc_RuntimeError,
>                 "Unexpected error in an Op's C code. "
>                 "No Python exception was set.");
>             }
>         goto __label_6;}
>         }
>         /* First check if we are the base type exactly (the most common case),
>            then do the full subclass check if needed. */
>         if (py_V5->ob_type != &PyGpuArrayType &&
>             !PyObject_TypeCheck(py_V5, &PyGpuArrayType)) {
>             PyErr_SetString(PyExc_ValueError, "expected a GpuArray");
>             {
>         __failure = 6;
>         if (!PyErr_Occurred()) {
>             PyErr_SetString(PyExc_RuntimeError,
>                 "Unexpected error in an Op's C code. "
>                 "No Python exception was set.");
>             }
>         goto __label_6;}
>         }
>         V5 = (PyGpuArrayObject *)py_V5;
>         Py_INCREF(V5);
>         
> {
> 
>     py_V7 = PyList_GET_ITEM(storage_V7, 0);
>     {Py_XINCREF(py_V7);}
>     
>         V7 = NULL;
>         if (py_V7 == Py_None) {
>             PyErr_SetString(PyExc_ValueError, "expected a GpuArray, not 
> None");
>             {
>         __failure = 8;
>         if (!PyErr_Occurred()) {
>             PyErr_SetString(PyExc_RuntimeError,
>                 "Unexpected error in an Op's C code. "
>                 "No Python exception was set.");
>             }
>         goto __label_8;}
>         }
>         /* First check if we are the base type exactly (the most common case),
>            then do the full subclass check if needed. */
>         if (py_V7->ob_type != &PyGpuArrayType &&
>             !PyObject_TypeCheck(py_V7, &PyGpuArrayType)) {
>             PyErr_SetString(PyExc_ValueError, "expected a GpuArray");
>             {
>         __failure = 8;
>         if (!PyErr_Occurred()) {
>             PyErr_SetString(PyExc_RuntimeError,
>                 "Unexpected error in an Op's C code. "
>                 "No Python exception was set.");
>             }
>         goto __label_8;}
>         }
>         V7 = (PyGpuArrayObject *)py_V7;
>         Py_INCREF(V7);
>         
> {
> 
>     py_V9 = PyList_GET_ITEM(storage_V9, 0);
>     {Py_XINCREF(py_V9);}
>     
>   V9 = (cudnnConvolutionDescriptor_t)PyCapsule_GetPointer(py_V9, NULL);
>   if (V9 == NULL) {
>         __failure = 10;
>         if (!PyErr_Occurred()) {
>             PyErr_SetString(PyExc_RuntimeError,
>                 "Unexpected error in an Op's C code. "
>                 "No Python exception was set.");
>             }
>         goto __label_10;}
>         
> {
> 
>     py_V11 = PyList_GET_ITEM(storage_V11, 0);
>     {Py_XINCREF(py_V11);}
>     
>             if (!PyObject_TypeCheck(py_V11, &PyFloat32ArrType_Type))
>             {
>                 PyErr_Format(PyExc_ValueError,
>                     "Scalar check failed (npy_float32)");
>                 {
>         __failure = 12;
>         if (!PyErr_Occurred()) {
>             PyErr_SetString(PyExc_RuntimeError,
>                 "Unexpected error in an Op's C code. "
>                 "No Python exception was set.");
>             }
>         goto __label_12;}
>             }
>             
>         PyArray_ScalarAsCtype(py_V11, &V11);
>         
> {
> 
>     py_V13 = PyList_GET_ITEM(storage_V13, 0);
>     {Py_XINCREF(py_V13);}
>     
>             if (!PyObject_TypeCheck(py_V13, &PyFloat32ArrType_Type))
>             {
>                 PyErr_Format(PyExc_ValueError,
>                     "Scalar check failed (npy_float32)");
>                 {
>         __failure = 14;
>         if (!PyErr_Occurred()) {
>             PyErr_SetString(PyExc_RuntimeError,
>                 "Unexpected error in an Op's C code. "
>                 "No Python exception was set.");
>             }
>         goto __label_14;}
>             }
>             
>         PyArray_ScalarAsCtype(py_V13, &V13);
>         
> {
> 
> {
> // Op class GpuDnnConv
> 
>                 #define APPLY_SPECIFIC(str) 
> str##_node_d7db26d1cc9fb59ee56289f4dda72577_0
> 
> #define CONV_INPLACE 1
> 
> #define CONV_ALGO CUDNN_CONVOLUTION_FWD_ALGO_IMPLICIT_PRECOMP_GEMM
>                 {
>                   if (APPLY_SPECIFIC(conv_fwd)(V3, V5, V7, V9, V11, V13, &V1, 
> V15) != 0) {
>                     {
>         __failure = 17;
>         if (!PyErr_Occurred()) {
>             PyErr_SetString(PyExc_RuntimeError,
>                 "Unexpected error in an Op's C code. "
>                 "No Python exception was set.");
>             }
>         goto __label_17;}
>                   }
>                 }
>                 #undef APPLY_SPECIFIC
> 
> #undef CONV_INPLACE
> 
> #undef CONV_ALGO
>                 __label_17:
> 
> double __DUMMY_17;
> 
> }
> __label_16:
> 
> double __DUMMY_16;
> 
> }
> __label_14:
> 
>     {Py_XDECREF(py_V13);}
>     
> double __DUMMY_14;
> 
> }
> __label_12:
> 
>     {Py_XDECREF(py_V11);}
>     
> double __DUMMY_12;
> 
> }
> __label_10:
> 
>     {Py_XDECREF(py_V9);}
>     
> double __DUMMY_10;
> 
> }
> __label_8:
> Py_XDECREF(V7); V7 = NULL;
>     {Py_XDECREF(py_V7);}
>     
> double __DUMMY_8;
> 
> }
> __label_6:
> Py_XDECREF(V5); V5 = NULL;
>     {Py_XDECREF(py_V5);}
>     
> double __DUMMY_6;
> 
> }
> __label_4:
> Py_XDECREF(V3); V3 = NULL;
>     {Py_XDECREF(py_V3);}
>     
> double __DUMMY_4;
> 
> }
> __label_2:
> 
>     if (!__failure) {
>       
>         if (!V1) {
>             Py_XDECREF(py_V1);
>             Py_INCREF(Py_None);
>             py_V1 = Py_None;
>         } else if ((void *)py_V1 != (void *)V1) {
>             Py_XDECREF(py_V1);
>             py_V1 = (PyObject *)V1;
>             Py_INCREF(py_V1);
>         }
>         
>       PyObject* old = PyList_GET_ITEM(storage_V1, 0);
>       {Py_XINCREF(py_V1);}
>       PyList_SET_ITEM(storage_V1, 0, py_V1);
>       {Py_XDECREF(old);}
>     }
>     Py_XDECREF(V1); V1 = NULL;
>     {Py_XDECREF(py_V1);}
>     
> double __DUMMY_2;
> 
> }
> 
>             
>         if (__failure) {
>             // When there is a failure, this code puts the exception
>             // in __ERROR.
>             PyObject* err_type = NULL;
>             PyObject* err_msg = NULL;
>             PyObject* err_traceback = NULL;
>             PyErr_Fetch(&err_type, &err_msg, &err_traceback);
>             if (!err_type) {err_type = Py_None;Py_INCREF(Py_None);}
>             if (!err_msg) {err_msg = Py_None; Py_INCREF(Py_None);}
>             if (!err_traceback) {err_traceback = Py_None; Py_INCREF(Py_None);}
>             PyObject* old_err_type = PyList_GET_ITEM(__ERROR, 0);
>             PyObject* old_err_msg = PyList_GET_ITEM(__ERROR, 1);
>             PyObject* old_err_traceback = PyList_GET_ITEM(__ERROR, 2);
>             PyList_SET_ITEM(__ERROR, 0, err_type);
>             PyList_SET_ITEM(__ERROR, 1, err_msg);
>             PyList_SET_ITEM(__ERROR, 2, err_traceback);
>             {Py_XDECREF(old_err_type);}
>             {Py_XDECREF(old_err_msg);}
>             {Py_XDECREF(old_err_traceback);}
>         }
>         // The failure code is returned to index what code block failed.
>         return __failure;
>         
>         }
>     };
>     }
>     
> 
>         static int 
> __struct_compiled_op_d7db26d1cc9fb59ee56289f4dda72577_executor(__struct_compiled_op_d7db26d1cc9fb59ee56289f4dda72577*
>  self) {
>             return self->run();
>         }
> 
>         static void 
> __struct_compiled_op_d7db26d1cc9fb59ee56289f4dda72577_destructor(void* 
> executor, void* self) {
>             delete 
> ((__struct_compiled_op_d7db26d1cc9fb59ee56289f4dda72577*)self);
>         }
>         
> //////////////////////
> ////  Functions
> //////////////////////
> static PyObject * instantiate(PyObject * self, PyObject *argtuple) {
>   assert(PyTuple_Check(argtuple));
>   if (9 != PyTuple_Size(argtuple)){ 
>      PyErr_Format(PyExc_TypeError, "Wrong number of arguments, expected 9, 
> got %i", (int)PyTuple_Size(argtuple));
>      return NULL;
>   }
>   __struct_compiled_op_d7db26d1cc9fb59ee56289f4dda72577* struct_ptr = new 
> __struct_compiled_op_d7db26d1cc9fb59ee56289f4dda72577();
>   if (struct_ptr->init( PyTuple_GET_ITEM(argtuple, 
> 0),PyTuple_GET_ITEM(argtuple, 1),PyTuple_GET_ITEM(argtuple, 
> 2),PyTuple_GET_ITEM(argtuple, 3),PyTuple_GET_ITEM(argtuple, 
> 4),PyTuple_GET_ITEM(argtuple, 5),PyTuple_GET_ITEM(argtuple, 
> 6),PyTuple_GET_ITEM(argtuple, 7),PyTuple_GET_ITEM(argtuple, 8) ) != 0) {
>     delete struct_ptr;
>     return NULL;
>   }
>   PyObject* thunk = 
> PyCObject_FromVoidPtrAndDesc((void*)(&__struct_compiled_op_d7db26d1cc9fb59ee56289f4dda72577_executor),
>  struct_ptr, 
> __struct_compiled_op_d7db26d1cc9fb59ee56289f4dda72577_destructor);
>   return thunk; }
> 
> //////////////////////
> ////  Module init
> //////////////////////
> static PyMethodDef MyMethods[] = {
>       {"instantiate", instantiate, METH_VARARGS, "undocumented"} ,
>       {NULL, NULL, 0, NULL}
> };
> PyMODINIT_FUNC initd7db26d1cc9fb59ee56289f4dda72577(void){
>    import_pygpu__gpuarray();
>    import_array();
>    
> 
> setup_ext_cuda();
> 
> 
>    (void) Py_InitModule("d7db26d1cc9fb59ee56289f4dda72577", MyMethods);
> }


-- 
Pascal

-- 

--- 
You received this message because you are subscribed to the Google Groups 
"theano-users" group.
To unsubscribe from this group and stop receiving emails from it, send an email 
to theano-users+unsubscr...@googlegroups.com.
For more options, visit https://groups.google.com/d/optout.

Reply via email to