Make _handle apply specific.

theano/sandbox/gpuarray/dnn_base.c

@@ -102,17 +102,19 @@ setup_ext_cuda();
 
 #section support_code_struct
 
-cudnnHandle_t _handle;
+cudnnHandle_t APPLY_SPECIFIC(_handle);
 
 #section init_code_struct
 
-cuda_enter(pygpu_default_context()->ctx);
-cudnnStatus_t err;
-_handle = NULL;
-if ((err = cudnnCreate(&_handle)) != CUDNN_STATUS_SUCCESS) {
-  PyErr_Format(PyExc_RuntimeError, "could not create cuDNN handle: %s",
-               cudnnGetErrorString(err));
+{
+  cuda_enter(pygpu_default_context()->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(pygpu_default_context()->ctx);
+    FAIL;
+  }
   cuda_exit(pygpu_default_context()->ctx);
-  FAIL;
 }
-cuda_exit(pygpu_default_context()->ctx);

theano/sandbox/gpuarray/dnn_fwd.c

@@ -80,7 +80,7 @@ APPLY_SPECIFIC(conv_fwd)(PyGpuArrayObject *input, PyGpuArrayObject *kerns,
     int count;
     cudnnConvolutionFwdAlgoPerf_t choice;
     err = cudnnFindConvolutionForwardAlgorithm(
-      _handle, APPLY_SPECIFIC(input), APPLY_SPECIFIC(kerns),
+      APPLY_SPECIFIC(_handle), APPLY_SPECIFIC(input), APPLY_SPECIFIC(kerns),
       desc, APPLY_SPECIFIC(output), 1, &count, &choice);
 
     if (err != CUDNN_STATUS_SUCCESS) {
@@ -103,7 +103,7 @@ APPLY_SPECIFIC(conv_fwd)(PyGpuArrayObject *input, PyGpuArrayObject *kerns,
     }
 
     err = cudnnGetConvolutionForwardAlgorithm(
-      _handle, APPLY_SPECIFIC(input), APPLY_SPECIFIC(kerns),
+      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) {
@@ -181,7 +181,7 @@ APPLY_SPECIFIC(conv_fwd)(PyGpuArrayObject *input, PyGpuArrayObject *kerns,
   {
     size_t worksize;
     gpudata *workspace;
-    err = cudnnGetConvolutionForwardWorkspaceSize(_handle,
+    err = cudnnGetConvolutionForwardWorkspaceSize(APPLY_SPECIFIC(_handle),
                                                   APPLY_SPECIFIC(input),
                                                   APPLY_SPECIFIC(kerns),
                                                   desc,
@@ -212,7 +212,7 @@ APPLY_SPECIFIC(conv_fwd)(PyGpuArrayObject *input, PyGpuArrayObject *kerns,
     }
 
     err = cudnnConvolutionForward(
-      _handle,
+      APPLY_SPECIFIC(_handle),
       alpha_p,
       APPLY_SPECIFIC(input), PyGpuArray_DEV_DATA(input),
       APPLY_SPECIFIC(kerns), PyGpuArray_DEV_DATA(kerns),

theano/sandbox/gpuarray/dnn_gi.c

@@ -79,7 +79,7 @@ APPLY_SPECIFIC(conv_gi)(PyGpuArrayObject *kerns, PyGpuArrayObject *output,
     cudnnConvolutionBwdDataAlgoPerf_t choice;
 
     err = cudnnFindConvolutionBackwardDataAlgorithm(
-      _handle, APPLY_SPECIFIC(input), APPLY_SPECIFIC(output), desc,
+      APPLY_SPECIFIC(_handle), APPLY_SPECIFIC(input), APPLY_SPECIFIC(output), desc,
       APPLY_SPECIFIC(kerns), 1, &count, &choice);
 
     if (err != CUDNN_STATUS_SUCCESS) {
@@ -102,7 +102,7 @@ APPLY_SPECIFIC(conv_gi)(PyGpuArrayObject *kerns, PyGpuArrayObject *output,
     }
 
     err = cudnnGetConvolutionBackwardDataAlgorithm(
-      _handle, APPLY_SPECIFIC(input), APPLY_SPECIFIC(output),
+      APPLY_SPECIFIC(_handle), APPLY_SPECIFIC(input), APPLY_SPECIFIC(output),
       desc, APPLY_SPECIFIC(kerns),
       CUDNN_CONVOLUTION_BWD_DATA_SPECIFY_WORKSPACE_LIMIT, free, &algo);
     if (err != CUDNN_STATUS_SUCCESS) {
@@ -157,7 +157,7 @@ APPLY_SPECIFIC(conv_gi)(PyGpuArrayObject *kerns, PyGpuArrayObject *output,
   gpudata *workspace;
 
   err = cudnnGetConvolutionBackwardDataWorkspaceSize(
-    _handle, APPLY_SPECIFIC(kerns), APPLY_SPECIFIC(output), desc,
+    APPLY_SPECIFIC(_handle), APPLY_SPECIFIC(kerns), APPLY_SPECIFIC(output), desc,
     APPLY_SPECIFIC(input), algo, &worksize);
 
   if (err != CUDNN_STATUS_SUCCESS) {
@@ -178,7 +178,7 @@ APPLY_SPECIFIC(conv_gi)(PyGpuArrayObject *kerns, PyGpuArrayObject *output,
   }
 
   err = cudnnConvolutionBackwardData_v3(
-    _handle,
+    APPLY_SPECIFIC(_handle),
     alpha_p,
     APPLY_SPECIFIC(kerns), PyGpuArray_DEV_DATA(kerns),
     APPLY_SPECIFIC(output), PyGpuArray_DEV_DATA(output),

theano/sandbox/gpuarray/dnn_gw.c

@@ -79,7 +79,7 @@ APPLY_SPECIFIC(conv_gw)(PyGpuArrayObject *input, PyGpuArrayObject *output,
     cudnnConvolutionBwdFilterAlgoPerf_t choice;
 
     err = cudnnFindConvolutionBackwardFilterAlgorithm(
-      _handle, APPLY_SPECIFIC(input), APPLY_SPECIFIC(output), desc,
+      APPLY_SPECIFIC(_handle), APPLY_SPECIFIC(input), APPLY_SPECIFIC(output), desc,
       APPLY_SPECIFIC(kerns), 1, &count, &choice);
 
     if (err != CUDNN_STATUS_SUCCESS) {
@@ -103,7 +103,7 @@ APPLY_SPECIFIC(conv_gw)(PyGpuArrayObject *input, PyGpuArrayObject *output,
     }
 
     err = cudnnGetConvolutionBackwardFilterAlgorithm(
-      _handle, APPLY_SPECIFIC(input), APPLY_SPECIFIC(output),
+      APPLY_SPECIFIC(_handle), APPLY_SPECIFIC(input), APPLY_SPECIFIC(output),
       desc, APPLY_SPECIFIC(kerns),
       CUDNN_CONVOLUTION_BWD_FILTER_SPECIFY_WORKSPACE_LIMIT, free, &algo);
     if (err != CUDNN_STATUS_SUCCESS) {
@@ -159,7 +159,7 @@ APPLY_SPECIFIC(conv_gw)(PyGpuArrayObject *input, PyGpuArrayObject *output,
   gpudata *workspace;
 
   err = cudnnGetConvolutionBackwardFilterWorkspaceSize(
-    _handle, APPLY_SPECIFIC(input), APPLY_SPECIFIC(output), desc,
+    APPLY_SPECIFIC(_handle), APPLY_SPECIFIC(input), APPLY_SPECIFIC(output), desc,
     APPLY_SPECIFIC(kerns), algo, &worksize);
 
   if (err != CUDNN_STATUS_SUCCESS) {
@@ -179,7 +179,7 @@ APPLY_SPECIFIC(conv_gw)(PyGpuArrayObject *input, PyGpuArrayObject *output,
   }
 
   err = cudnnConvolutionBackwardFilter_v3(
-    _handle,
+    APPLY_SPECIFIC(_handle),
     alpha_p,
     APPLY_SPECIFIC(input), PyGpuArray_DEV_DATA(input),
     APPLY_SPECIFIC(output), PyGpuArray_DEV_DATA(output),

theano/sandbox/gpuarray/dnn_pool.c

@@ -76,7 +76,7 @@ int APPLY_SPECIFIC(dnn_pool)(PyGpuArrayObject *img,
 
     cuda_enter(c->ctx);
     err = cudnnPoolingForward(
-      _handle, desc,
+      APPLY_SPECIFIC(_handle), desc,
       &alpha,
       APPLY_SPECIFIC(input), PyGpuArray_DEV_DATA(img),
       &beta,

theano/sandbox/gpuarray/dnn_pool_grad.c

@@ -94,7 +94,7 @@ int APPLY_SPECIFIC(dnn_pool_grad)(PyGpuArrayObject *inp,
 
     cuda_enter(c->ctx);
     err = cudnnPoolingBackward(
-      _handle, desc,
+      APPLY_SPECIFIC(_handle), desc,
       &alpha,
       APPLY_SPECIFIC(output), PyGpuArray_DEV_DATA(out),
       APPLY_SPECIFIC(output_grad), PyGpuArray_DEV_DATA(out_grad),

theano/sandbox/gpuarray/dnn_softmax.c

@@ -55,7 +55,7 @@ int APPLY_SPECIFIC(softmax)(PyGpuArrayObject *x,
 
     cuda_enter(c->ctx);
     err = cudnnSoftmaxForward(
-      _handle,
+      APPLY_SPECIFIC(_handle),
       SOFTMAX_ALGO,
       SOFTMAX_MODE,
       (void *)&alpha,

theano/sandbox/gpuarray/dnn_softmax_grad.c

@@ -68,7 +68,7 @@ int APPLY_SPECIFIC(softmax_grad)(PyGpuArrayObject *dy,
 
     cuda_enter(c->ctx);
     err = cudnnSoftmaxBackward(
-      _handle,
+      APPLY_SPECIFIC(_handle),
       SOFTMAX_ALGO,
       SOFTMAX_MODE,
       (void *)&alpha,