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+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 *)α
> > beta_p = (void *)β
> > break;
> > case GA_FLOAT:
> > case GA_HALF:
> > alpha_p = (void *)⁡
> > 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.
--
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.
