diff --git a/src/conv_lstm_layer.c b/src/conv_lstm_layer.c index 6cbaf1c3..5da2bab3 100644 --- a/src/conv_lstm_layer.c +++ b/src/conv_lstm_layer.c @@ -66,44 +66,44 @@ layer make_conv_lstm_layer(int batch, int h, int w, int c, int output_filters, i // U l.uf = (layer*)calloc(1, sizeof(layer)); - *(l.uf) = make_convolutional_layer(batch, steps, h, w, c, output_filters, groups, size, stride, dilation, pad, activation, batch_normalize, 0, xnor, 0, 0, 0, NULL); + *(l.uf) = make_convolutional_layer(batch, steps, h, w, c, output_filters, groups, size, stride, stride, dilation, pad, activation, batch_normalize, 0, xnor, 0, 0, 0, NULL); l.uf->batch = batch; if (l.workspace_size < l.uf->workspace_size) l.workspace_size = l.uf->workspace_size; l.ui = (layer*)calloc(1, sizeof(layer)); - *(l.ui) = make_convolutional_layer(batch, steps, h, w, c, output_filters, groups, size, stride, dilation, pad, activation, batch_normalize, 0, xnor, 0, 0, 0, NULL); + *(l.ui) = make_convolutional_layer(batch, steps, h, w, c, output_filters, groups, size, stride, stride, dilation, pad, activation, batch_normalize, 0, xnor, 0, 0, 0, NULL); l.ui->batch = batch; if (l.workspace_size < l.ui->workspace_size) l.workspace_size = l.ui->workspace_size; l.ug = (layer*)calloc(1, sizeof(layer)); - *(l.ug) = make_convolutional_layer(batch, steps, h, w, c, output_filters, groups, size, stride, dilation, pad, activation, batch_normalize, 0, xnor, 0, 0, 0, NULL); + *(l.ug) = make_convolutional_layer(batch, steps, h, w, c, output_filters, groups, size, stride, stride, dilation, pad, activation, batch_normalize, 0, xnor, 0, 0, 0, NULL); l.ug->batch = batch; if (l.workspace_size < l.ug->workspace_size) l.workspace_size = l.ug->workspace_size; l.uo = (layer*)calloc(1, sizeof(layer)); - *(l.uo) = make_convolutional_layer(batch, steps, h, w, c, output_filters, groups, size, stride, dilation, pad, activation, batch_normalize, 0, xnor, 0, 0, 0, NULL); + *(l.uo) = make_convolutional_layer(batch, steps, h, w, c, output_filters, groups, size, stride, stride, dilation, pad, activation, batch_normalize, 0, xnor, 0, 0, 0, NULL); l.uo->batch = batch; if (l.workspace_size < l.uo->workspace_size) l.workspace_size = l.uo->workspace_size; // W l.wf = (layer*)calloc(1, sizeof(layer)); - *(l.wf) = make_convolutional_layer(batch, steps, h, w, output_filters, output_filters, groups, size, stride, dilation, pad, activation, batch_normalize, 0, xnor, 0, 0, 0, NULL); + *(l.wf) = make_convolutional_layer(batch, steps, h, w, output_filters, output_filters, groups, size, stride, stride, dilation, pad, activation, batch_normalize, 0, xnor, 0, 0, 0, NULL); l.wf->batch = batch; if (l.workspace_size < l.wf->workspace_size) l.workspace_size = l.wf->workspace_size; l.wi = (layer*)calloc(1, sizeof(layer)); - *(l.wi) = make_convolutional_layer(batch, steps, h, w, output_filters, output_filters, groups, size, stride, dilation, pad, activation, batch_normalize, 0, xnor, 0, 0, 0, NULL); + *(l.wi) = make_convolutional_layer(batch, steps, h, w, output_filters, output_filters, groups, size, stride, stride, dilation, pad, activation, batch_normalize, 0, xnor, 0, 0, 0, NULL); l.wi->batch = batch; if (l.workspace_size < l.wi->workspace_size) l.workspace_size = l.wi->workspace_size; l.wg = (layer*)calloc(1, sizeof(layer)); - *(l.wg) = make_convolutional_layer(batch, steps, h, w, output_filters, output_filters, groups, size, stride, dilation, pad, activation, batch_normalize, 0, xnor, 0, 0, 0, NULL); + *(l.wg) = make_convolutional_layer(batch, steps, h, w, output_filters, output_filters, groups, size, stride, stride, dilation, pad, activation, batch_normalize, 0, xnor, 0, 0, 0, NULL); l.wg->batch = batch; if (l.workspace_size < l.wg->workspace_size) l.workspace_size = l.wg->workspace_size; l.wo = (layer*)calloc(1, sizeof(layer)); - *(l.wo) = make_convolutional_layer(batch, steps, h, w, output_filters, output_filters, groups, size, stride, dilation, pad, activation, batch_normalize, 0, xnor, 0, 0, 0, NULL); + *(l.wo) = make_convolutional_layer(batch, steps, h, w, output_filters, output_filters, groups, size, stride, stride, dilation, pad, activation, batch_normalize, 0, xnor, 0, 0, 0, NULL); l.wo->batch = batch; if (l.workspace_size < l.wo->workspace_size) l.workspace_size = l.wo->workspace_size; @@ -111,21 +111,21 @@ layer make_conv_lstm_layer(int batch, int h, int w, int c, int output_filters, i // V l.vf = (layer*)calloc(1, sizeof(layer)); if (l.peephole) { - *(l.vf) = make_convolutional_layer(batch, steps, h, w, output_filters, output_filters, groups, size, stride, dilation, pad, activation, batch_normalize, 0, xnor, 0, 0, 0, NULL); + *(l.vf) = make_convolutional_layer(batch, steps, h, w, output_filters, output_filters, groups, size, stride, stride, dilation, pad, activation, batch_normalize, 0, xnor, 0, 0, 0, NULL); l.vf->batch = batch; if (l.workspace_size < l.vf->workspace_size) l.workspace_size = l.vf->workspace_size; } l.vi = (layer*)calloc(1, sizeof(layer)); if (l.peephole) { - *(l.vi) = make_convolutional_layer(batch, steps, h, w, output_filters, output_filters, groups, size, stride, dilation, pad, activation, batch_normalize, 0, xnor, 0, 0, 0, NULL); + *(l.vi) = make_convolutional_layer(batch, steps, h, w, output_filters, output_filters, groups, size, stride, stride, dilation, pad, activation, batch_normalize, 0, xnor, 0, 0, 0, NULL); l.vi->batch = batch; if (l.workspace_size < l.vi->workspace_size) l.workspace_size = l.vi->workspace_size; } l.vo = (layer*)calloc(1, sizeof(layer)); if (l.peephole) { - *(l.vo) = make_convolutional_layer(batch, steps, h, w, output_filters, output_filters, groups, size, stride, dilation, pad, activation, batch_normalize, 0, xnor, 0, 0, 0, NULL); + *(l.vo) = make_convolutional_layer(batch, steps, h, w, output_filters, output_filters, groups, size, stride, stride, dilation, pad, activation, batch_normalize, 0, xnor, 0, 0, 0, NULL); l.vo->batch = batch; if (l.workspace_size < l.vo->workspace_size) l.workspace_size = l.vo->workspace_size; } diff --git a/src/convolutional_kernels.cu b/src/convolutional_kernels.cu index e404ecab..07a0a0d7 100644 --- a/src/convolutional_kernels.cu +++ b/src/convolutional_kernels.cu @@ -177,7 +177,7 @@ void forward_convolutional_layer_gpu(convolutional_layer l, network_state state) fast_binarize_weights_gpu(l.weights_gpu, l.n, (l.c / l.groups)*l.size*l.size, l.binary_weights_gpu, l.mean_arr_gpu); } - if (l.align_bit_weights_gpu && !state.train && l.c >= 32) + if (l.align_bit_weights_gpu && !state.train && l.c >= 32 && l.stride_x == l.stride_y) { //return; //cudaError_t status = cudaSuccess; @@ -574,7 +574,7 @@ void forward_convolutional_layer_gpu(convolutional_layer l, network_state state) l.h, l.w, // input size (h, w) l.size, l.size, // kernel size (h, w) l.pad, l.pad, // padding (h, w) - l.stride, l.stride, // stride (h, w) + l.stride_y, l.stride_x, // stride (h, w) l.dilation, l.dilation, // dilation (h, w) state.workspace); // output @@ -819,7 +819,7 @@ void backward_convolutional_layer_gpu(convolutional_layer l, network_state state l.h, l.w, // input size (h, w) l.size, l.size, // kernel size (h, w) l.pad, l.pad, // padding (h, w) - l.stride, l.stride, // stride (h, w) + l.stride_y, l.stride_x, // stride (h, w) l.dilation, l.dilation, // dilation (h, w) state.workspace); // output //gemm_ongpu(0, 1, m, n, k, 1, a + i*m*k, k, b, k, 1, c, n); @@ -844,7 +844,7 @@ void backward_convolutional_layer_gpu(convolutional_layer l, network_state state l.h, l.w, // input size (h, w) l.size, l.size, // kernel size (h, w) l.pad, l.pad, // padding size (h, w) - l.stride, l.stride, // stride size (h, w) + l.stride_y, l.stride_x, // stride size (h, w) l.dilation, l.dilation, // dilation size (h, w) delta); // output (delta) diff --git a/src/convolutional_layer.c b/src/convolutional_layer.c index 93ac79a0..207e3f27 100644 --- a/src/convolutional_layer.c +++ b/src/convolutional_layer.c @@ -76,12 +76,12 @@ void binarize_input(float *input, int n, int size, float *binary) int convolutional_out_height(convolutional_layer l) { - return (l.h + 2*l.pad - l.size) / l.stride + 1; + return (l.h + 2*l.pad - l.size) / l.stride_y + 1; } int convolutional_out_width(convolutional_layer l) { - return (l.w + 2*l.pad - l.size) / l.stride + 1; + return (l.w + 2*l.pad - l.size) / l.stride_x + 1; } image get_convolutional_image(convolutional_layer l) @@ -276,9 +276,9 @@ void cudnn_convolutional_setup(layer *l, int cudnn_preference) //printf("\n l->dilation = %d, l->pad = %d, l->size = %d \n", l->dilation, l->pad, l->size); #if(CUDNN_MAJOR >= 6) - CHECK_CUDNN(cudnnSetConvolution2dDescriptor(l->convDesc, l->pad * l->dilation, l->pad* l->dilation, l->stride, l->stride, l->dilation, l->dilation, CUDNN_CROSS_CORRELATION, CUDNN_DATA_FLOAT)); // cudnn >= 6.0 + CHECK_CUDNN(cudnnSetConvolution2dDescriptor(l->convDesc, l->pad * l->dilation, l->pad* l->dilation, l->stride_y, l->stride_x, l->dilation, l->dilation, CUDNN_CROSS_CORRELATION, CUDNN_DATA_FLOAT)); // cudnn >= 6.0 #else - CHECK_CUDNN(cudnnSetConvolution2dDescriptor(l->convDesc, l->pad * l->dilation, l->pad * l->dilation, l->stride, l->stride, l->dilation, l->dilation, CUDNN_CROSS_CORRELATION)); // cudnn 5.1 + CHECK_CUDNN(cudnnSetConvolution2dDescriptor(l->convDesc, l->pad * l->dilation, l->pad * l->dilation, l->stride_y, l->stride_x, l->dilation, l->dilation, CUDNN_CROSS_CORRELATION)); // cudnn 5.1 #endif int forward_algo = CUDNN_CONVOLUTION_FWD_PREFER_FASTEST; int backward_algo = CUDNN_CONVOLUTION_BWD_DATA_PREFER_FASTEST; @@ -332,7 +332,7 @@ void cudnn_convolutional_setup(layer *l, int cudnn_preference) #endif #endif -convolutional_layer make_convolutional_layer(int batch, int steps, int h, int w, int c, int n, int groups, int size, int stride, int dilation, int padding, ACTIVATION activation, int batch_normalize, int binary, int xnor, int adam, int use_bin_output, int index, convolutional_layer *share_layer) +convolutional_layer make_convolutional_layer(int batch, int steps, int h, int w, int c, int n, int groups, int size, int stride_x, int stride_y, int dilation, int padding, ACTIVATION activation, int batch_normalize, int binary, int xnor, int adam, int use_bin_output, int index, convolutional_layer *share_layer) { int total_batch = batch*steps; int i; @@ -354,7 +354,9 @@ convolutional_layer make_convolutional_layer(int batch, int steps, int h, int w, l.use_bin_output = use_bin_output; l.batch = batch; l.steps = steps; - l.stride = stride; + l.stride = stride_x; + l.stride_y = stride_x; + l.stride_x = stride_y; l.dilation = dilation; l.size = size; l.pad = padding; @@ -553,11 +555,14 @@ convolutional_layer make_convolutional_layer(int batch, int steps, int h, int w, else if(l.share_layer) fprintf(stderr, "convS "); else fprintf(stderr, "conv "); - if(groups > 1) fprintf(stderr, "%5d/%4d ", n, groups); + if (groups > 1) fprintf(stderr, "%5d/%4d ", n, groups); else fprintf(stderr, "%5d ", n); - if(dilation > 1) fprintf(stderr, "%2d x%2d/%2d(%1d)", size, size, stride, dilation); - else fprintf(stderr, "%2d x%2d/%2d ", size, size, stride); + if (stride_x != stride_y) fprintf(stderr, "%2d x%2d/%2dx%2d ", size, size, stride_x, stride_y); + else { + if (dilation > 1) fprintf(stderr, "%2d x%2d/%2d(%1d)", size, size, stride_x, dilation); + else fprintf(stderr, "%2d x%2d/%2d ", size, size, stride_x); + } fprintf(stderr, "%4d x%4d x%4d -> %4d x%4d x%4d %5.3f BF\n", w, h, c, l.out_w, l.out_h, l.out_c, l.bflops); @@ -583,7 +588,7 @@ void denormalize_convolutional_layer(convolutional_layer l) void test_convolutional_layer() { - convolutional_layer l = make_convolutional_layer(1, 1, 5, 5, 3, 2, 1, 5, 2, 1, 1, LEAKY, 1, 0, 0, 0, 0, 0, NULL); + convolutional_layer l = make_convolutional_layer(1, 1, 5, 5, 3, 2, 1, 5, 2, 2, 1, 1, LEAKY, 1, 0, 0, 0, 0, 0, NULL); l.batch_normalize = 1; float data[] = {1,1,1,1,1, 1,1,1,1,1, @@ -921,7 +926,7 @@ void forward_convolutional_layer(convolutional_layer l, network_state state) //gemm(0,0,m,n,k,1,a,k,b,n,1,c,n); //gemm_nn_custom(m, n, k, 1, a, k, b, n, c, n); - if (l.xnor && l.align_bit_weights && !state.train) + if (l.xnor && l.align_bit_weights && !state.train && l.stride_x == l.stride_y) { memset(b, 0, l.bit_align*l.size*l.size*l.c * sizeof(float)); @@ -1053,7 +1058,7 @@ void forward_convolutional_layer(convolutional_layer l, network_state state) l.h, l.w, // input size (h, w) l.size, l.size, // kernel size (h, w) l.pad, l.pad, // padding (h, w) - l.stride, l.stride, // stride (h, w) + l.stride_y, l.stride_x, // stride (h, w) l.dilation, l.dilation, // dilation (h, w) b); // output @@ -1229,7 +1234,7 @@ void backward_convolutional_layer(convolutional_layer l, network_state state) l.h, l.w, // input size (h, w) l.size, l.size, // kernel size (h, w) l.pad, l.pad, // padding (h, w) - l.stride, l.stride, // stride (h, w) + l.stride_y, l.stride_x, // stride (h, w) l.dilation, l.dilation, // dilation (h, w) b); // output @@ -1251,7 +1256,7 @@ void backward_convolutional_layer(convolutional_layer l, network_state state) l.h, l.w, // input size (h, w) l.size, l.size, // kernel size (h, w) l.pad, l.pad, // padding (h, w) - l.stride, l.stride, // stride (h, w) + l.stride_y, l.stride_x, // stride (h, w) l.dilation, l.dilation, // dilation (h, w) state.delta + (i*l.groups + j)* (l.c / l.groups)*l.h*l.w); // output (delta) } diff --git a/src/convolutional_layer.h b/src/convolutional_layer.h index e62b155c..1167175c 100644 --- a/src/convolutional_layer.h +++ b/src/convolutional_layer.h @@ -30,7 +30,7 @@ void cuda_convert_f32_to_f16(float* input_f32, size_t size, float *output_f16); #endif size_t get_convolutional_workspace_size(layer l); -convolutional_layer make_convolutional_layer(int batch, int steps, int h, int w, int c, int n, int groups, int size, int stride, int dilation, int padding, ACTIVATION activation, int batch_normalize, int binary, int xnor, int adam, int use_bin_output, int index, convolutional_layer *share_layer); +convolutional_layer make_convolutional_layer(int batch, int steps, int h, int w, int c, int n, int groups, int size, int stride_x, int stride_y, int dilation, int padding, ACTIVATION activation, int batch_normalize, int binary, int xnor, int adam, int use_bin_output, int index, convolutional_layer *share_layer); void denormalize_convolutional_layer(convolutional_layer l); void resize_convolutional_layer(convolutional_layer *layer, int w, int h); void forward_convolutional_layer(const convolutional_layer layer, network_state state); diff --git a/src/crnn_layer.c b/src/crnn_layer.c index 7609003b..eaded279 100644 --- a/src/crnn_layer.c +++ b/src/crnn_layer.c @@ -50,17 +50,17 @@ layer make_crnn_layer(int batch, int h, int w, int c, int hidden_filters, int ou l.state = (float*)calloc(l.hidden * l.batch * (l.steps + 1), sizeof(float)); l.input_layer = (layer*)calloc(1, sizeof(layer)); - *(l.input_layer) = make_convolutional_layer(batch, steps, h, w, c, hidden_filters, groups, size, stride, dilation, pad, activation, batch_normalize, 0, xnor, 0, 0, 0, NULL); + *(l.input_layer) = make_convolutional_layer(batch, steps, h, w, c, hidden_filters, groups, size, stride, stride, dilation, pad, activation, batch_normalize, 0, xnor, 0, 0, 0, NULL); l.input_layer->batch = batch; if (l.workspace_size < l.input_layer->workspace_size) l.workspace_size = l.input_layer->workspace_size; l.self_layer = (layer*)calloc(1, sizeof(layer)); - *(l.self_layer) = make_convolutional_layer(batch, steps, h, w, hidden_filters, hidden_filters, groups, size, stride, dilation, pad, activation, batch_normalize, 0, xnor, 0, 0, 0, NULL); + *(l.self_layer) = make_convolutional_layer(batch, steps, h, w, hidden_filters, hidden_filters, groups, size, stride, stride, dilation, pad, activation, batch_normalize, 0, xnor, 0, 0, 0, NULL); l.self_layer->batch = batch; if (l.workspace_size < l.self_layer->workspace_size) l.workspace_size = l.self_layer->workspace_size; l.output_layer = (layer*)calloc(1, sizeof(layer)); - *(l.output_layer) = make_convolutional_layer(batch, steps, h, w, hidden_filters, output_filters, groups, size, stride, dilation, pad, activation, batch_normalize, 0, xnor, 0, 0, 0, NULL); + *(l.output_layer) = make_convolutional_layer(batch, steps, h, w, hidden_filters, output_filters, groups, size, stride, stride, dilation, pad, activation, batch_normalize, 0, xnor, 0, 0, 0, NULL); l.output_layer->batch = batch; if (l.workspace_size < l.output_layer->workspace_size) l.workspace_size = l.output_layer->workspace_size; diff --git a/src/parser.c b/src/parser.c index 48bd42bd..8283f7ed 100644 --- a/src/parser.c +++ b/src/parser.c @@ -158,6 +158,8 @@ convolutional_layer parse_convolutional(list *options, size_params params, netwo int groups = option_find_int_quiet(options, "groups", 1); int size = option_find_int(options, "size",1); int stride = option_find_int(options, "stride",1); + int stride_x = option_find_int_quiet(options, "stride_x", stride); + int stride_y = option_find_int_quiet(options, "stride_y", stride); int dilation = option_find_int_quiet(options, "dilation", 1); if (size == 1) dilation = 1; int pad = option_find_int_quiet(options, "pad",0); @@ -167,9 +169,10 @@ convolutional_layer parse_convolutional(list *options, size_params params, netwo char *activation_s = option_find_str(options, "activation", "logistic"); ACTIVATION activation = get_activation(activation_s); - int share_index = option_find_int_quiet(options, "share_index", -1); + int share_index = option_find_int_quiet(options, "share_index", -1000000000); convolutional_layer *share_layer = NULL; - if(share_index > -1) share_layer = &net.layers[share_index]; + if(share_index >= 0) share_layer = &net.layers[share_index]; + else if(share_index != -1000000000) share_layer = &net.layers[params.index + share_index]; int batch,h,w,c; h = params.h; @@ -182,7 +185,7 @@ convolutional_layer parse_convolutional(list *options, size_params params, netwo int xnor = option_find_int_quiet(options, "xnor", 0); int use_bin_output = option_find_int_quiet(options, "bin_output", 0); - convolutional_layer layer = make_convolutional_layer(batch,1,h,w,c,n,groups,size,stride,dilation,padding,activation, batch_normalize, binary, xnor, params.net.adam, use_bin_output, params.index, share_layer); + convolutional_layer layer = make_convolutional_layer(batch,1,h,w,c,n,groups,size,stride_x,stride_y,dilation,padding,activation, batch_normalize, binary, xnor, params.net.adam, use_bin_output, params.index, share_layer); layer.flipped = option_find_int_quiet(options, "flipped", 0); layer.dot = option_find_float_quiet(options, "dot", 0); layer.assisted_excitation = option_find_float_quiet(options, "assisted_excitation", 0);