Added support for Tensor Cores CC >= 7.0 (V100). For FP16/32 (mixed precisi...

AlexeyAB

2018-02-25 cad4d1618fee74471d335314cb77070fee951a42

 src/network.c

@@ -50,6 +50,7 @@
    int batch_num = get_current_batch(net);
    int i;
    float rate;
   if (batch_num < net.burn_in) return net.learning_rate * pow((float)batch_num / net.burn_in, net.power);
    switch (net.policy) {
        case CONSTANT:
            return net.learning_rate;
@@ -66,8 +67,9 @@
        case EXP:
            return net.learning_rate * pow(net.gamma, batch_num);
        case POLY:
            if (batch_num < net.burn_in) return net.learning_rate * pow((float)batch_num / net.burn_in, net.power);
            return net.learning_rate * pow(1 - (float)batch_num / net.max_batches, net.power);
         return net.learning_rate * pow(1 - (float)batch_num / net.max_batches, net.power);
            //if (batch_num < net.burn_in) return net.learning_rate * pow((float)batch_num / net.burn_in, net.power);
            //return net.learning_rate * pow(1 - (float)batch_num / net.max_batches, net.power);
        case RANDOM:
            return net.learning_rate * pow(rand_uniform(0,1), net.power);
        case SIG:
@@ -218,6 +220,7 @@
            state.delta = prev.delta;
        }
        layer l = net.layers[i];
        if (l.stopbackward) break;
        l.backward(l, state);
    }
}
@@ -313,7 +316,20 @@
        net->layers[i].batch = b;
#ifdef CUDNN
        if(net->layers[i].type == CONVOLUTIONAL){
            cudnn_convolutional_setup(net->layers + i);
         cudnn_convolutional_setup(net->layers + i, cudnn_fastest);
         /*
         layer *l = net->layers + i;
            cudnn_convolutional_setup(l, cudnn_fastest);
         // check for excessive memory consumption 
         size_t free_byte;
         size_t total_byte;
         check_error(cudaMemGetInfo(&free_byte, &total_byte));
         if (l->workspace_size > free_byte || l->workspace_size >= total_byte / 2) {
            printf(" used slow CUDNN algo without Workspace! \n");
            cudnn_convolutional_setup(l, cudnn_smallest);
            l->workspace_size = get_workspace_size(*l);
         }
         */
        }
#endif
    }
@@ -325,6 +341,12 @@
    cuda_set_device(net->gpu_index);
    if(gpu_index >= 0){
        cuda_free(net->workspace);
      if (net->input_gpu) {
         cuda_free(*net->input_gpu);
         *net->input_gpu = 0;
         cuda_free(*net->truth_gpu);
         *net->truth_gpu = 0;
      }
    }
#endif
    int i;
@@ -337,6 +359,7 @@
    //fflush(stderr);
    for (i = 0; i < net->n; ++i){
        layer l = net->layers[i];
      //printf(" %d: layer = %d,", i, l.type);
        if(l.type == CONVOLUTIONAL){
            resize_convolutional_layer(&l, w, h);
        }else if(l.type == CROP){
@@ -356,6 +379,7 @@
        }else if(l.type == COST){
            resize_cost_layer(&l, inputs);
        }else{
         fprintf(stderr, "Resizing type %d \n", (int)l.type);
            error("Cannot resize this type of layer");
        }
        if(l.workspace_size > workspace_size) workspace_size = l.workspace_size;
@@ -367,13 +391,9 @@
    }
#ifdef GPU
    if(gpu_index >= 0){
        if(net->input_gpu) {
            cuda_free(*net->input_gpu);
            *net->input_gpu = 0;
            cuda_free(*net->truth_gpu);
            *net->truth_gpu = 0;
        }
      printf(" try to allocate workspace = %zu * sizeof(float), ", (workspace_size - 1) / sizeof(float) + 1);
        net->workspace = cuda_make_array(0, (workspace_size-1)/sizeof(float)+1);
      printf(" CUDA allocate done! \n");
    }else {
        free(net->workspace);
        net->workspace = calloc(1, workspace_size);
@@ -590,15 +610,19 @@

void free_network(network net)
{
    int i;
    for(i = 0; i < net.n; ++i){
        free_layer(net.layers[i]);
    }
    free(net.layers);
   int i;
   for (i = 0; i < net.n; ++i) {
      free_layer(net.layers[i]);
   }
   free(net.layers);
#ifdef GPU
    if(*net.input_gpu) cuda_free(*net.input_gpu);
    if(*net.truth_gpu) cuda_free(*net.truth_gpu);
    if(net.input_gpu) free(net.input_gpu);
    if(net.truth_gpu) free(net.truth_gpu);
   if (gpu_index >= 0) cuda_free(net.workspace);
   else free(net.workspace);
   if (*net.input_gpu) cuda_free(*net.input_gpu);
   if (*net.truth_gpu) cuda_free(*net.truth_gpu);
   if (net.input_gpu) free(net.input_gpu);
   if (net.truth_gpu) free(net.truth_gpu);
#else
   free(net.workspace);
#endif
}