From a6b2511a566f77a0838dc1dd0d5f3e3c49a8faa0 Mon Sep 17 00:00:00 2001
From: Joseph Redmon <pjreddie@gmail.com>
Date: Sat, 25 Jun 2016 23:13:54 +0000
Subject: [PATCH] idk
---
src/convolutional_layer.c | 128 +++++++++++++++++++-----------------------
1 files changed, 58 insertions(+), 70 deletions(-)
diff --git a/src/convolutional_layer.c b/src/convolutional_layer.c
index f0c312c..4014a24 100644
--- a/src/convolutional_layer.c
+++ b/src/convolutional_layer.c
@@ -8,8 +8,13 @@
#include <stdio.h>
#include <time.h>
+#ifdef AI2
+#include "xnor_layer.h"
+#endif
+
#ifndef AI2
#define AI2 0
+void forward_xnor_layer(layer l, network_state state);
#endif
void swap_binary(convolutional_layer *l)
@@ -40,6 +45,14 @@
}
}
+void binarize_cpu(float *input, int n, float *binary)
+{
+ int i;
+ for(i = 0; i < n; ++i){
+ binary[i] = (input[i] > 0) ? 1 : -1;
+ }
+}
+
void binarize_input(float *input, int n, int size, float *binary)
{
int i, s;
@@ -123,6 +136,47 @@
#endif
}
+#ifdef GPU
+#ifdef CUDNN
+void cudnn_convolutional_setup(layer *l)
+{
+ cudnnSetTensor4dDescriptor(l->dsrcTensorDesc, CUDNN_TENSOR_NCHW, CUDNN_DATA_FLOAT, l->batch, l->c, l->h, l->w);
+ cudnnSetTensor4dDescriptor(l->ddstTensorDesc, CUDNN_TENSOR_NCHW, CUDNN_DATA_FLOAT, l->batch, l->out_c, l->out_h, l->out_w);
+ cudnnSetFilter4dDescriptor(l->dfilterDesc, CUDNN_DATA_FLOAT, CUDNN_TENSOR_NCHW, l->n, l->c, l->size, l->size);
+
+ cudnnSetTensor4dDescriptor(l->srcTensorDesc, CUDNN_TENSOR_NCHW, CUDNN_DATA_FLOAT, l->batch, l->c, l->h, l->w);
+ cudnnSetTensor4dDescriptor(l->dstTensorDesc, CUDNN_TENSOR_NCHW, CUDNN_DATA_FLOAT, l->batch, l->out_c, l->out_h, l->out_w);
+ cudnnSetFilter4dDescriptor(l->filterDesc, CUDNN_DATA_FLOAT, CUDNN_TENSOR_NCHW, l->n, l->c, l->size, l->size);
+ int padding = l->pad ? l->size/2 : 0;
+ cudnnSetConvolution2dDescriptor(l->convDesc, padding, padding, l->stride, l->stride, 1, 1, CUDNN_CROSS_CORRELATION);
+ cudnnGetConvolutionForwardAlgorithm(cudnn_handle(),
+ l->srcTensorDesc,
+ l->filterDesc,
+ l->convDesc,
+ l->dstTensorDesc,
+ CUDNN_CONVOLUTION_FWD_PREFER_FASTEST,
+ 0,
+ &l->fw_algo);
+ cudnnGetConvolutionBackwardDataAlgorithm(cudnn_handle(),
+ l->filterDesc,
+ l->ddstTensorDesc,
+ l->convDesc,
+ l->dsrcTensorDesc,
+ CUDNN_CONVOLUTION_BWD_DATA_PREFER_FASTEST,
+ 0,
+ &l->bd_algo);
+ cudnnGetConvolutionBackwardFilterAlgorithm(cudnn_handle(),
+ l->srcTensorDesc,
+ l->ddstTensorDesc,
+ l->convDesc,
+ l->dfilterDesc,
+ CUDNN_CONVOLUTION_BWD_FILTER_PREFER_FASTEST,
+ 0,
+ &l->bf_algo);
+}
+#endif
+#endif
+
convolutional_layer make_convolutional_layer(int batch, int h, int w, int c, int n, int size, int stride, int pad, ACTIVATION activation, int batch_normalize, int binary, int xnor)
{
int i;
@@ -227,39 +281,7 @@
cudnnCreateTensorDescriptor(&l.ddstTensorDesc);
cudnnCreateFilterDescriptor(&l.dfilterDesc);
cudnnCreateConvolutionDescriptor(&l.convDesc);
- cudnnSetTensor4dDescriptor(l.dsrcTensorDesc, CUDNN_TENSOR_NCHW, CUDNN_DATA_FLOAT, l.batch, l.c, l.h, l.w);
- cudnnSetTensor4dDescriptor(l.ddstTensorDesc, CUDNN_TENSOR_NCHW, CUDNN_DATA_FLOAT, l.batch, l.out_c, l.out_h, l.out_w);
- cudnnSetFilter4dDescriptor(l.dfilterDesc, CUDNN_DATA_FLOAT, CUDNN_TENSOR_NCHW, l.n, l.c, l.size, l.size);
-
- cudnnSetTensor4dDescriptor(l.srcTensorDesc, CUDNN_TENSOR_NCHW, CUDNN_DATA_FLOAT, l.batch, l.c, l.h, l.w);
- cudnnSetTensor4dDescriptor(l.dstTensorDesc, CUDNN_TENSOR_NCHW, CUDNN_DATA_FLOAT, l.batch, l.out_c, l.out_h, l.out_w);
- cudnnSetFilter4dDescriptor(l.filterDesc, CUDNN_DATA_FLOAT, CUDNN_TENSOR_NCHW, l.n, l.c, l.size, l.size);
- int padding = l.pad ? l.size/2 : 0;
- cudnnSetConvolution2dDescriptor(l.convDesc, padding, padding, l.stride, l.stride, 1, 1, CUDNN_CROSS_CORRELATION);
- cudnnGetConvolutionForwardAlgorithm(cudnn_handle(),
- l.srcTensorDesc,
- l.filterDesc,
- l.convDesc,
- l.dstTensorDesc,
- CUDNN_CONVOLUTION_FWD_PREFER_FASTEST,
- 0,
- &l.fw_algo);
- cudnnGetConvolutionBackwardDataAlgorithm(cudnn_handle(),
- l.filterDesc,
- l.ddstTensorDesc,
- l.convDesc,
- l.dsrcTensorDesc,
- CUDNN_CONVOLUTION_BWD_DATA_PREFER_FASTEST,
- 0,
- &l.bd_algo);
- cudnnGetConvolutionBackwardFilterAlgorithm(cudnn_handle(),
- l.srcTensorDesc,
- l.ddstTensorDesc,
- l.convDesc,
- l.dfilterDesc,
- CUDNN_CONVOLUTION_BWD_FILTER_PREFER_FASTEST,
- 0,
- &l.bf_algo);
+ cudnn_convolutional_setup(&l);
#endif
#endif
l.workspace_size = get_workspace_size(l);
@@ -331,39 +353,7 @@
l->delta_gpu = cuda_make_array(l->delta, l->batch*out_h*out_w*l->n);
l->output_gpu = cuda_make_array(l->output, l->batch*out_h*out_w*l->n);
#ifdef CUDNN
- cudnnSetTensor4dDescriptor(l->dsrcTensorDesc, CUDNN_TENSOR_NCHW, CUDNN_DATA_FLOAT, l->batch, l->c, l->h, l->w);
- cudnnSetTensor4dDescriptor(l->ddstTensorDesc, CUDNN_TENSOR_NCHW, CUDNN_DATA_FLOAT, l->batch, l->out_c, l->out_h, l->out_w);
- cudnnSetFilter4dDescriptor(l->dfilterDesc, CUDNN_DATA_FLOAT, CUDNN_TENSOR_NCHW, l->n, l->c, l->size, l->size);
-
- cudnnSetTensor4dDescriptor(l->srcTensorDesc, CUDNN_TENSOR_NCHW, CUDNN_DATA_FLOAT, l->batch, l->c, l->h, l->w);
- cudnnSetTensor4dDescriptor(l->dstTensorDesc, CUDNN_TENSOR_NCHW, CUDNN_DATA_FLOAT, l->batch, l->out_c, l->out_h, l->out_w);
- cudnnSetFilter4dDescriptor(l->filterDesc, CUDNN_DATA_FLOAT, CUDNN_TENSOR_NCHW, l->n, l->c, l->size, l->size);
- int padding = l->pad ? l->size/2 : 0;
- cudnnSetConvolution2dDescriptor(l->convDesc, padding, padding, l->stride, l->stride, 1, 1, CUDNN_CROSS_CORRELATION);
- cudnnGetConvolutionForwardAlgorithm(cudnn_handle(),
- l->srcTensorDesc,
- l->filterDesc,
- l->convDesc,
- l->dstTensorDesc,
- CUDNN_CONVOLUTION_FWD_PREFER_FASTEST,
- 0,
- &l->fw_algo);
- cudnnGetConvolutionBackwardDataAlgorithm(cudnn_handle(),
- l->filterDesc,
- l->ddstTensorDesc,
- l->convDesc,
- l->dsrcTensorDesc,
- CUDNN_CONVOLUTION_BWD_DATA_PREFER_FASTEST,
- 0,
- &l->bd_algo);
- cudnnGetConvolutionBackwardFilterAlgorithm(cudnn_handle(),
- l->srcTensorDesc,
- l->ddstTensorDesc,
- l->convDesc,
- l->dfilterDesc,
- CUDNN_CONVOLUTION_BWD_FILTER_PREFER_FASTEST,
- 0,
- &l->bf_algo);
+ cudnn_convolutional_setup(l);
#endif
#endif
l->workspace_size = get_workspace_size(*l);
@@ -444,12 +434,10 @@
}
*/
- if(l.xnor && (l.c%32 != 0 || !AI2)){
+ if(l.xnor ){
binarize_filters(l.filters, l.n, l.c*l.size*l.size, l.binary_filters);
swap_binary(&l);
- for(i = 0; i < l.batch; ++i){
- binarize_input(state.input + i*l.inputs, l.c, l.h*l.w, l.binary_input + i*l.inputs);
- }
+ binarize_cpu(state.input, l.c*l.h*l.w*l.batch, l.binary_input);
state.input = l.binary_input;
}
--
Gitblit v1.10.0