| | |
| | | return float_to_image(w,h,c,l.delta); |
| | | } |
| | | |
| | | size_t get_workspace_size(layer l){ |
| | | #ifdef CUDNN |
| | | size_t most = 0; |
| | | size_t s = 0; |
| | | cudnnGetConvolutionForwardWorkspaceSize(cudnn_handle(), |
| | | l.srcTensorDesc, |
| | | l.filterDesc, |
| | | l.convDesc, |
| | | l.dstTensorDesc, |
| | | l.fw_algo, |
| | | &s); |
| | | if (s > most) most = s; |
| | | cudnnGetConvolutionBackwardFilterWorkspaceSize(cudnn_handle(), |
| | | l.srcTensorDesc, |
| | | l.ddstTensorDesc, |
| | | l.convDesc, |
| | | l.dfilterDesc, |
| | | l.bf_algo, |
| | | &s); |
| | | if (s > most) most = s; |
| | | cudnnGetConvolutionBackwardDataWorkspaceSize(cudnn_handle(), |
| | | l.filterDesc, |
| | | l.ddstTensorDesc, |
| | | l.convDesc, |
| | | l.dsrcTensorDesc, |
| | | l.bd_algo, |
| | | &s); |
| | | if (s > most) most = s; |
| | | return most; |
| | | #else |
| | | return (size_t)l.out_h*l.out_w*l.size*l.size*l.c*sizeof(float); |
| | | #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; |
| | |
| | | l.outputs = l.out_h * l.out_w * l.out_c; |
| | | l.inputs = l.w * l.h * l.c; |
| | | |
| | | l.col_image = calloc(out_h*out_w*size*size*c, sizeof(float)); |
| | | l.output = calloc(l.batch*out_h * out_w * n, sizeof(float)); |
| | | l.delta = calloc(l.batch*out_h * out_w * n, sizeof(float)); |
| | | |
| | |
| | | l.scales_gpu = cuda_make_array(l.scales, n); |
| | | l.scale_updates_gpu = cuda_make_array(l.scale_updates, n); |
| | | |
| | | l.col_image_gpu = cuda_make_array(l.col_image, out_h*out_w*size*size*c); |
| | | l.delta_gpu = cuda_make_array(l.delta, l.batch*out_h*out_w*n); |
| | | l.output_gpu = cuda_make_array(l.output, l.batch*out_h*out_w*n); |
| | | |
| | |
| | | l.x_gpu = cuda_make_array(l.output, l.batch*out_h*out_w*n); |
| | | l.x_norm_gpu = cuda_make_array(l.output, l.batch*out_h*out_w*n); |
| | | } |
| | | #ifdef CUDNN |
| | | cudnnCreateTensorDescriptor(&l.srcTensorDesc); |
| | | cudnnCreateTensorDescriptor(&l.dstTensorDesc); |
| | | cudnnCreateFilterDescriptor(&l.filterDesc); |
| | | cudnnCreateTensorDescriptor(&l.dsrcTensorDesc); |
| | | 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); |
| | | #endif |
| | | #endif |
| | | l.workspace_size = get_workspace_size(l); |
| | | l.activation = activation; |
| | | |
| | | fprintf(stderr, "Convolutional Layer: %d x %d x %d image, %d filters -> %d x %d x %d image\n", h,w,c,n, out_h, out_w, n); |
| | |
| | | l->outputs = l->out_h * l->out_w * l->out_c; |
| | | l->inputs = l->w * l->h * l->c; |
| | | |
| | | l->col_image = realloc(l->col_image, |
| | | out_h*out_w*l->size*l->size*l->c*sizeof(float)); |
| | | l->output = realloc(l->output, |
| | | l->batch*out_h * out_w * l->n*sizeof(float)); |
| | | l->delta = realloc(l->delta, |
| | | l->batch*out_h * out_w * l->n*sizeof(float)); |
| | | |
| | | #ifdef GPU |
| | | cuda_free(l->col_image_gpu); |
| | | cuda_free(l->delta_gpu); |
| | | cuda_free(l->output_gpu); |
| | | |
| | | l->col_image_gpu = cuda_make_array(l->col_image, out_h*out_w*l->size*l->size*l->c); |
| | | 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); |
| | | #endif |
| | | #endif |
| | | l->workspace_size = get_workspace_size(*l); |
| | | } |
| | | |
| | | void add_bias(float *output, float *biases, int batch, int n, int size) |
| | |
| | | |
| | | fill_cpu(l.outputs*l.batch, 0, l.output, 1); |
| | | /* |
| | | if(l.binary){ |
| | | binarize_filters(l.filters, l.n, l.c*l.size*l.size, l.binary_filters); |
| | | binarize_filters2(l.filters, l.n, l.c*l.size*l.size, l.cfilters, l.scales); |
| | | swap_binary(&l); |
| | | } |
| | | */ |
| | | if(l.binary){ |
| | | binarize_filters(l.filters, l.n, l.c*l.size*l.size, l.binary_filters); |
| | | binarize_filters2(l.filters, l.n, l.c*l.size*l.size, l.cfilters, l.scales); |
| | | swap_binary(&l); |
| | | } |
| | | */ |
| | | |
| | | if(l.binary){ |
| | | int m = l.n; |
| | |
| | | int n = out_h*out_w; |
| | | |
| | | char *a = l.cfilters; |
| | | float *b = l.col_image; |
| | | float *b = state.workspace; |
| | | float *c = l.output; |
| | | |
| | | for(i = 0; i < l.batch; ++i){ |
| | |
| | | int n = out_h*out_w; |
| | | |
| | | float *a = l.filters; |
| | | float *b = l.col_image; |
| | | float *b = state.workspace; |
| | | float *c = l.output; |
| | | |
| | | for(i = 0; i < l.batch; ++i){ |
| | |
| | | |
| | | for(i = 0; i < l.batch; ++i){ |
| | | float *a = l.delta + i*m*k; |
| | | float *b = l.col_image; |
| | | float *b = state.workspace; |
| | | float *c = l.filter_updates; |
| | | |
| | | float *im = state.input+i*l.c*l.h*l.w; |
| | |
| | | if(state.delta){ |
| | | a = l.filters; |
| | | b = l.delta + i*m*k; |
| | | c = l.col_image; |
| | | c = state.workspace; |
| | | |
| | | gemm(1,0,n,k,m,1,a,n,b,k,0,c,k); |
| | | |
| | | col2im_cpu(l.col_image, l.c, l.h, l.w, l.size, l.stride, l.pad, state.delta+i*l.c*l.h*l.w); |
| | | col2im_cpu(state.workspace, l.c, l.h, l.w, l.size, l.stride, l.pad, state.delta+i*l.c*l.h*l.w); |
| | | } |
| | | } |
| | | } |