AlexeyAB
2018-03-30 23cb35e6c8eae8b59fab161036ae3f417a55c8db
src/network_kernels.cu
@@ -51,6 +51,8 @@
            fill_ongpu(l.outputs * l.batch, 0, l.delta_gpu, 1);
        }
        l.forward_gpu(l, state);
      if(net.wait_stream)
         cudaStreamSynchronize(get_cuda_stream());
        state.input = l.output_gpu;
    }
}
@@ -64,6 +66,7 @@
    for(i = net.n-1; i >= 0; --i){
        state.index = i;
        layer l = net.layers[i];
        if (l.stopbackward) break;
        if(i == 0){
            state.input = original_input;
            state.delta = original_delta;
@@ -78,11 +81,13 @@
void update_network_gpu(network net)
{
    cuda_set_device(net.gpu_index);
    int i;
    int update_batch = net.batch*net.subdivisions;
    float rate = get_current_rate(net);
    for(i = 0; i < net.n; ++i){
        layer l = net.layers[i];
        l.t = get_current_batch(net);
        if(l.update_gpu){
            l.update_gpu(l, update_batch, rate, net.momentum, net.decay);
        }
@@ -108,7 +113,15 @@
    state.delta = 0;
    state.truth = *net.truth_gpu;
    state.train = 1;
#ifdef CUDNN_HALF
   int i;
   for (i = 0; i < net.n; ++i) {
      layer l = net.layers[i];
      cuda_convert_f32_to_f16(l.weights_gpu, l.c*l.n*l.size*l.size, l.weights_gpu16);
   }
#endif
    forward_network_gpu(net, state);
   //cudaStreamSynchronize(get_cuda_stream());
    backward_network_gpu(net, state);
}
@@ -134,7 +147,6 @@
    free(ptr);
    cuda_set_device(args.net.gpu_index);
    *args.err = train_network(args.net, args.d);
    printf("%d\n", args.net.gpu_index);
    return 0;
}
@@ -177,6 +189,7 @@
{
    int update_batch = net.batch*net.subdivisions;
    float rate = get_current_rate(net);
    l.t = get_current_batch(net);
    if(l.update_gpu){
        l.update_gpu(l, update_batch, rate, net.momentum, net.decay);
    }
@@ -376,7 +389,7 @@
float *get_network_output_layer_gpu(network net, int i)
{
    layer l = net.layers[i];
    cuda_pull_array(l.output_gpu, l.output, l.outputs*l.batch);
    if(l.type != REGION) cuda_pull_array(l.output_gpu, l.output, l.outputs*l.batch);
    return l.output;
}
@@ -389,7 +402,8 @@
float *network_predict_gpu(network net, float *input)
{
    cuda_set_device(net.gpu_index);
   if (net.gpu_index != cuda_get_device())
      cuda_set_device(net.gpu_index);
    int size = get_network_input_size(net) * net.batch;
    network_state state;
    state.index = 0;