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 [email protected].
> 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 [email protected].
For more options, visit https://groups.google.com/d/optout.
