I'm using the master branch version (9?). The compiler is Visual Studio 
2013. The file I posted was created from an unmodified library. Perhaps 
it's a unix vs. windows line ending problem? I'll try to remove the extra 
lines to see if it compiles.

Strangely, during the error dump there are no extra lines:
00598   #define CONV_ALGO CUDNN_CONVOLUTION_FWD_ALGO_IMPLICIT_PRECOMP_GEMM
00599   #define FAIL { \
00600           if (!PyErr_Occurred()) { \
00601               PyErr_SetString(PyExc_RuntimeError, \
00602                   "Unexpected error in an Op's C code. " \
00603                   "No Python exception was set."); \
00604               } \
00605           return 17; \
00606   }
00607   #define PARAMS V15
00608
00609
00610   {
00611     // We need to keep a reference here to have it available in the 
destructor.
00612     ctx = PARAMS;
00613     Py_INCREF(ctx);



On Thursday, September 22, 2016 at 7:47:55 PM UTC+1, Pascal Lamblin wrote:
>
> Also, an issue seems to be that there are blank lines interspersed 
> through the definition of the FAIL macro, which prevents its correct 
> definition (and give the parse errors). 
>
> I have no idea why this is the case, though. 
> Which version of Theano are you using? 
> Have you modified the source in any way? 
> Can you try with 0.8.2? 
>
> On Thu, Sep 22, 2016, Pascal Lamblin wrote: 
> > 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...@googlegroups.com <javascript:>. 
> > > 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...@googlegroups.com <javascript:>. 
> > For more options, visit https://groups.google.com/d/optout. 
>
> -- 
> 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