void forward_batchnorm_layer_gpu(layer l, network_state state) { if(l.type == BATCHNORM) copy_ongpu(l.outputs*l.batch, state.input, 1, l.output_gpu, 1); if(l.type == CONNECTED){ l.out_c = l.outputs; l.out_h = l.out_w = 1; } if (state.train) { fast_mean_gpu(l.output_gpu, l.batch, l.out_c, l.out_h*l.out_w, l.mean_gpu); fast_variance_gpu(l.output_gpu, l.mean_gpu, l.batch, l.out_c, l.out_h*l.out_w, l.variance_gpu); scal_ongpu(l.out_c, .99, l.rolling_mean_gpu, 1); axpy_ongpu(l.out_c, .01, l.mean_gpu, 1, l.rolling_mean_gpu, 1); scal_ongpu(l.out_c, .99, l.rolling_variance_gpu, 1); axpy_ongpu(l.out_c, .01, l.variance_gpu, 1, l.rolling_variance_gpu, 1); copy_ongpu(l.outputs*l.batch, l.output_gpu, 1, l.x_gpu, 1); normalize_gpu(l.output_gpu, l.mean_gpu, l.variance_gpu, l.batch, l.out_c, l.out_h*l.out_w); copy_ongpu(l.outputs*l.batch, l.output_gpu, 1, l.x_norm_gpu, 1); } else { normalize_gpu(l.output_gpu, l.rolling_mean_gpu, l.rolling_variance_gpu, l.batch, l.out_c, l.out_h*l.out_w); } scale_bias_gpu(l.output_gpu, l.scales_gpu, l.batch, l.out_c, l.out_h*l.out_w); }
void update_connected_layer_gpu(connected_layer l, int batch, float learning_rate, float momentum, float decay) { axpy_ongpu(l.outputs, learning_rate/batch, l.bias_updates_gpu, 1, l.biases_gpu, 1); scal_ongpu(l.outputs, momentum, l.bias_updates_gpu, 1); axpy_ongpu(l.inputs*l.outputs, -decay*batch, l.weights_gpu, 1, l.weight_updates_gpu, 1); axpy_ongpu(l.inputs*l.outputs, learning_rate/batch, l.weight_updates_gpu, 1, l.weights_gpu, 1); scal_ongpu(l.inputs*l.outputs, momentum, l.weight_updates_gpu, 1); }
void update_local_layer_gpu(local_layer l, int batch, float learning_rate, float momentum, float decay, cudaStream_t stream) { int locations = l.out_w*l.out_h; int size = l.size*l.size*l.c*l.n*locations; axpy_ongpu(l.outputs, learning_rate/batch, l.bias_updates_gpu, 1, l.biases_gpu, 1, stream); scal_ongpu(l.outputs, momentum, l.bias_updates_gpu, 1, stream); axpy_ongpu(size, -decay*batch, l.weights_gpu, 1, l.weight_updates_gpu, 1, stream); axpy_ongpu(size, learning_rate/batch, l.weight_updates_gpu, 1, l.weights_gpu, 1, stream); scal_ongpu(size, momentum, l.weight_updates_gpu, 1, stream); }
void update_local_layer_gpu(local_layer l, int batch, float learning_rate, float momentum, float decay) { int locations = l.out_w*l.out_h; int size = l.size*l.size*l.c*l.n*locations; axpy_ongpu(l.outputs, learning_rate/batch, l.bias_updates_gpu, 1, l.biases_gpu, 1); scal_ongpu(l.outputs, momentum, l.bias_updates_gpu, 1); axpy_ongpu(size, -decay*batch, l.filters_gpu, 1, l.filter_updates_gpu, 1); axpy_ongpu(size, learning_rate/batch, l.filter_updates_gpu, 1, l.filters_gpu, 1); scal_ongpu(size, momentum, l.filter_updates_gpu, 1); }
void forward_batchnorm_layer_gpu(layer l, network_state state) { if(l.type == BATCHNORM) copy_ongpu(l.outputs*l.batch, state.input, 1, l.output_gpu, 1); if(l.type == CONNECTED){ l.out_c = l.outputs; l.out_h = l.out_w = 1; } if (state.train) { #ifdef CUDNN copy_ongpu(l.outputs*l.batch, l.output_gpu, 1, l.x_gpu, 1); float one = 1; float zero = 0; cudnnBatchNormalizationForwardTraining(cudnn_handle(), CUDNN_BATCHNORM_SPATIAL, &one, &zero, l.dstTensorDesc, l.x_gpu, l.dstTensorDesc, l.output_gpu, l.normTensorDesc, l.scales_gpu, l.biases_gpu, .01, l.rolling_mean_gpu, l.rolling_variance_gpu, .00001, l.mean_gpu, l.variance_gpu); #else fast_mean_gpu(l.output_gpu, l.batch, l.out_c, l.out_h*l.out_w, l.mean_gpu); fast_variance_gpu(l.output_gpu, l.mean_gpu, l.batch, l.out_c, l.out_h*l.out_w, l.variance_gpu); scal_ongpu(l.out_c, .99, l.rolling_mean_gpu, 1); axpy_ongpu(l.out_c, .01, l.mean_gpu, 1, l.rolling_mean_gpu, 1); scal_ongpu(l.out_c, .99, l.rolling_variance_gpu, 1); axpy_ongpu(l.out_c, .01, l.variance_gpu, 1, l.rolling_variance_gpu, 1); copy_ongpu(l.outputs*l.batch, l.output_gpu, 1, l.x_gpu, 1); normalize_gpu(l.output_gpu, l.mean_gpu, l.variance_gpu, l.batch, l.out_c, l.out_h*l.out_w); copy_ongpu(l.outputs*l.batch, l.output_gpu, 1, l.x_norm_gpu, 1); scale_bias_gpu(l.output_gpu, l.scales_gpu, l.batch, l.out_c, l.out_h*l.out_w); add_bias_gpu(l.output_gpu, l.biases_gpu, l.batch, l.out_c, l.out_w*l.out_h); #endif } else { normalize_gpu(l.output_gpu, l.rolling_mean_gpu, l.rolling_variance_gpu, l.batch, l.out_c, l.out_h*l.out_w); scale_bias_gpu(l.output_gpu, l.scales_gpu, l.batch, l.out_c, l.out_h*l.out_w); add_bias_gpu(l.output_gpu, l.biases_gpu, l.batch, l.out_c, l.out_w*l.out_h); } }
void forward_cost_layer_gpu(cost_layer l, network_state state) { if (!state.truth) return; if(l.smooth){ scal_ongpu(l.batch*l.inputs, (1-l.smooth), state.truth, 1); add_ongpu(l.batch*l.inputs, l.smooth * 1./l.inputs, state.truth, 1); } if (l.cost_type == MASKED) { mask_ongpu(l.batch*l.inputs, state.input, SECRET_NUM, state.truth); } if(l.cost_type == SMOOTH){ smooth_l1_gpu(l.batch*l.inputs, state.input, state.truth, l.delta_gpu, l.output_gpu); } else if (l.cost_type == L1){ l1_gpu(l.batch*l.inputs, state.input, state.truth, l.delta_gpu, l.output_gpu); } else { l2_gpu(l.batch*l.inputs, state.input, state.truth, l.delta_gpu, l.output_gpu); } if(l.ratio){ cuda_pull_array(l.delta_gpu, l.delta, l.batch*l.inputs); qsort(l.delta, l.batch*l.inputs, sizeof(float), float_abs_compare); int n = (1-l.ratio) * l.batch*l.inputs; float thresh = l.delta[n]; thresh = 0; printf("%f\n", thresh); supp_ongpu(l.batch*l.inputs, thresh, l.delta_gpu, 1); } if(l.thresh){ supp_ongpu(l.batch*l.inputs, l.thresh*1./l.inputs, l.delta_gpu, 1); } cuda_pull_array(l.output_gpu, l.output, l.batch*l.inputs); l.cost[0] = sum_array(l.output, l.batch*l.inputs); }
void train_lsd2(char *cfgfile, char *weightfile, char *acfgfile, char *aweightfile, int clear) { #ifdef GPU char *train_images = "/home/pjreddie/data/coco/trainvalno5k.txt"; char *backup_directory = "/home/pjreddie/backup/"; srand(time(0)); char *base = basecfg(cfgfile); printf("%s\n", base); network net = parse_network_cfg(cfgfile); if(weightfile){ load_weights(&net, weightfile); } if(clear) *net.seen = 0; char *abase = basecfg(acfgfile); network anet = parse_network_cfg(acfgfile); if(aweightfile){ load_weights(&anet, aweightfile); } if(clear) *anet.seen = 0; int i, j, k; layer imlayer = {}; for (i = 0; i < net.n; ++i) { if (net.layers[i].out_c == 3) { imlayer = net.layers[i]; break; } } printf("Learning Rate: %g, Momentum: %g, Decay: %g\n", net.learning_rate, net.momentum, net.decay); int imgs = net.batch*net.subdivisions; i = *net.seen/imgs; data train, buffer; list *plist = get_paths(train_images); //int N = plist->size; char **paths = (char **)list_to_array(plist); load_args args = {}; args.w = net.w; args.h = net.h; args.paths = paths; args.n = imgs; args.m = plist->size; args.d = &buffer; args.min = net.min_crop; args.max = net.max_crop; args.angle = net.angle; args.aspect = net.aspect; args.exposure = net.exposure; args.saturation = net.saturation; args.hue = net.hue; args.size = net.w; args.type = CLASSIFICATION_DATA; args.classes = 1; char *ls[1] = {"coco"}; args.labels = ls; pthread_t load_thread = load_data_in_thread(args); clock_t time; network_state gstate = {}; gstate.index = 0; gstate.net = net; int x_size = get_network_input_size(net)*net.batch; int y_size = 1*net.batch; gstate.input = cuda_make_array(0, x_size); gstate.truth = 0; gstate.delta = 0; gstate.train = 1; float *X = (float*)calloc(x_size, sizeof(float)); float *y = (float*)calloc(y_size, sizeof(float)); network_state astate = {}; astate.index = 0; astate.net = anet; int ay_size = get_network_output_size(anet)*anet.batch; astate.input = 0; astate.truth = 0; astate.delta = 0; astate.train = 1; float *imerror = cuda_make_array(0, imlayer.outputs); float *ones_gpu = cuda_make_array(0, ay_size); fill_ongpu(ay_size, 1, ones_gpu, 1); float aloss_avg = -1; float gloss_avg = -1; //data generated = copy_data(train); while (get_current_batch(net) < net.max_batches) { i += 1; time=clock(); pthread_join(load_thread, 0); train = buffer; load_thread = load_data_in_thread(args); printf("Loaded: %lf seconds\n", sec(clock()-time)); data generated = copy_data(train); time=clock(); float gloss = 0; for(j = 0; j < net.subdivisions; ++j){ get_next_batch(train, net.batch, j*net.batch, X, y); cuda_push_array(gstate.input, X, x_size); *net.seen += net.batch; forward_network_gpu(net, gstate); fill_ongpu(imlayer.outputs, 0, imerror, 1); astate.input = imlayer.output_gpu; astate.delta = imerror; astate.truth = ones_gpu; forward_network_gpu(anet, astate); backward_network_gpu(anet, astate); scal_ongpu(imlayer.outputs, 1, imerror, 1); axpy_ongpu(imlayer.outputs, 1, imerror, 1, imlayer.delta_gpu, 1); backward_network_gpu(net, gstate); printf("features %f\n", cuda_mag_array(imlayer.delta_gpu, imlayer.outputs)); printf("realness %f\n", cuda_mag_array(imerror, imlayer.outputs)); gloss += get_network_cost(net) /(net.subdivisions*net.batch); cuda_pull_array(imlayer.output_gpu, imlayer.output, x_size); for(k = 0; k < net.batch; ++k){ int index = j*net.batch + k; copy_cpu(imlayer.outputs, imlayer.output + k*imlayer.outputs, 1, generated.X.vals[index], 1); generated.y.vals[index][0] = 0; } } harmless_update_network_gpu(anet); data merge = concat_data(train, generated); randomize_data(merge); float aloss = train_network(anet, merge); update_network_gpu(net); update_network_gpu(anet); free_data(merge); free_data(train); free_data(generated); if (aloss_avg < 0) aloss_avg = aloss; aloss_avg = aloss_avg*.9 + aloss*.1; gloss_avg = gloss_avg*.9 + gloss*.1; printf("%d: gen: %f, adv: %f | gen_avg: %f, adv_avg: %f, %f rate, %lf seconds, %d images\n", i, gloss, aloss, gloss_avg, aloss_avg, get_current_rate(net), sec(clock()-time), i*imgs); if(i%1000==0){ char buff[256]; sprintf(buff, "%s/%s_%d.weights", backup_directory, base, i); save_weights(net, buff); sprintf(buff, "%s/%s_%d.weights", backup_directory, abase, i); save_weights(anet, buff); } if(i%100==0){ char buff[256]; sprintf(buff, "%s/%s.backup", backup_directory, base); save_weights(net, buff); sprintf(buff, "%s/%s.backup", backup_directory, abase); save_weights(anet, buff); } } char buff[256]; sprintf(buff, "%s/%s_final.weights", backup_directory, base); save_weights(net, buff); #endif }
void train_lsd3(char *fcfg, char *fweight, char *gcfg, char *gweight, char *acfg, char *aweight, int clear) { #ifdef GPU //char *train_images = "/home/pjreddie/data/coco/trainvalno5k.txt"; char *train_images = "/home/pjreddie/data/imagenet/imagenet1k.train.list"; //char *style_images = "/home/pjreddie/data/coco/trainvalno5k.txt"; char *style_images = "/home/pjreddie/zelda.txt"; char *backup_directory = "/home/pjreddie/backup/"; srand(time(0)); network fnet = load_network(fcfg, fweight, clear); network gnet = load_network(gcfg, gweight, clear); network anet = load_network(acfg, aweight, clear); char *gbase = basecfg(gcfg); char *abase = basecfg(acfg); printf("Learning Rate: %g, Momentum: %g, Decay: %g\n", gnet.learning_rate, gnet.momentum, gnet.decay); int imgs = gnet.batch*gnet.subdivisions; int i = *gnet.seen/imgs; data train, tbuffer; data style, sbuffer; list *slist = get_paths(style_images); char **spaths = (char **)list_to_array(slist); list *tlist = get_paths(train_images); char **tpaths = (char **)list_to_array(tlist); load_args targs= get_base_args(gnet); targs.paths = tpaths; targs.n = imgs; targs.m = tlist->size; targs.d = &tbuffer; targs.type = CLASSIFICATION_DATA; targs.classes = 1; char *ls[1] = {"zelda"}; targs.labels = ls; load_args sargs = get_base_args(gnet); sargs.paths = spaths; sargs.n = imgs; sargs.m = slist->size; sargs.d = &sbuffer; sargs.type = CLASSIFICATION_DATA; sargs.classes = 1; sargs.labels = ls; pthread_t tload_thread = load_data_in_thread(targs); pthread_t sload_thread = load_data_in_thread(sargs); clock_t time; float aloss_avg = -1; float floss_avg = -1; network_state fstate = {}; fstate.index = 0; fstate.net = fnet; int x_size = get_network_input_size(fnet)*fnet.batch; int y_size = get_network_output_size(fnet)*fnet.batch; fstate.input = cuda_make_array(0, x_size); fstate.truth = cuda_make_array(0, y_size); fstate.delta = cuda_make_array(0, x_size); fstate.train = 1; float *X = (float*)calloc(x_size, sizeof(float)); float *y = (float*)calloc(y_size, sizeof(float)); float *ones = cuda_make_array(0, anet.batch); float *zeros = cuda_make_array(0, anet.batch); fill_ongpu(anet.batch, .99, ones, 1); fill_ongpu(anet.batch, .01, zeros, 1); network_state astate = {}; astate.index = 0; astate.net = anet; int ax_size = get_network_input_size(anet)*anet.batch; int ay_size = get_network_output_size(anet)*anet.batch; astate.input = 0; astate.truth = ones; astate.delta = cuda_make_array(0, ax_size); astate.train = 1; network_state gstate = {}; gstate.index = 0; gstate.net = gnet; int gx_size = get_network_input_size(gnet)*gnet.batch; int gy_size = get_network_output_size(gnet)*gnet.batch; gstate.input = cuda_make_array(0, gx_size); gstate.truth = 0; gstate.delta = 0; gstate.train = 1; while (get_current_batch(gnet) < gnet.max_batches) { i += 1; time=clock(); pthread_join(tload_thread, 0); pthread_join(sload_thread, 0); train = tbuffer; style = sbuffer; tload_thread = load_data_in_thread(targs); sload_thread = load_data_in_thread(sargs); printf("Loaded: %lf seconds\n", sec(clock()-time)); data generated = copy_data(train); time=clock(); int j, k; float floss = 0; for(j = 0; j < fnet.subdivisions; ++j){ layer imlayer = gnet.layers[gnet.n - 1]; get_next_batch(train, fnet.batch, j*fnet.batch, X, y); cuda_push_array(fstate.input, X, x_size); cuda_push_array(gstate.input, X, gx_size); *gnet.seen += gnet.batch; forward_network_gpu(fnet, fstate); float *feats = fnet.layers[fnet.n - 2].output_gpu; copy_ongpu(y_size, feats, 1, fstate.truth, 1); forward_network_gpu(gnet, gstate); float *gen = gnet.layers[gnet.n-1].output_gpu; copy_ongpu(x_size, gen, 1, fstate.input, 1); fill_ongpu(x_size, 0, fstate.delta, 1); forward_network_gpu(fnet, fstate); backward_network_gpu(fnet, fstate); //HERE astate.input = gen; fill_ongpu(ax_size, 0, astate.delta, 1); forward_network_gpu(anet, astate); backward_network_gpu(anet, astate); float *delta = imlayer.delta_gpu; fill_ongpu(x_size, 0, delta, 1); scal_ongpu(x_size, 100, astate.delta, 1); scal_ongpu(x_size, .00001, fstate.delta, 1); axpy_ongpu(x_size, 1, fstate.delta, 1, delta, 1); axpy_ongpu(x_size, 1, astate.delta, 1, delta, 1); //fill_ongpu(x_size, 0, delta, 1); //cuda_push_array(delta, X, x_size); //axpy_ongpu(x_size, -1, imlayer.output_gpu, 1, delta, 1); //printf("pix error: %f\n", cuda_mag_array(delta, x_size)); printf("fea error: %f\n", cuda_mag_array(fstate.delta, x_size)); printf("adv error: %f\n", cuda_mag_array(astate.delta, x_size)); //axpy_ongpu(x_size, 1, astate.delta, 1, delta, 1); backward_network_gpu(gnet, gstate); floss += get_network_cost(fnet) /(fnet.subdivisions*fnet.batch); cuda_pull_array(imlayer.output_gpu, imlayer.output, x_size); for(k = 0; k < gnet.batch; ++k){ int index = j*gnet.batch + k; copy_cpu(imlayer.outputs, imlayer.output + k*imlayer.outputs, 1, generated.X.vals[index], 1); generated.y.vals[index][0] = .01; } } /* image sim = float_to_image(anet.w, anet.h, anet.c, style.X.vals[j]); show_image(sim, "style"); cvWaitKey(0); */ harmless_update_network_gpu(anet); data merge = concat_data(style, generated); randomize_data(merge); float aloss = train_network(anet, merge); update_network_gpu(gnet); free_data(merge); free_data(train); free_data(generated); free_data(style); if (aloss_avg < 0) aloss_avg = aloss; if (floss_avg < 0) floss_avg = floss; aloss_avg = aloss_avg*.9 + aloss*.1; floss_avg = floss_avg*.9 + floss*.1; printf("%d: gen: %f, adv: %f | gen_avg: %f, adv_avg: %f, %f rate, %lf seconds, %d images\n", i, floss, aloss, floss_avg, aloss_avg, get_current_rate(gnet), sec(clock()-time), i*imgs); if(i%1000==0){ char buff[256]; sprintf(buff, "%s/%s_%d.weights", backup_directory, gbase, i); save_weights(gnet, buff); sprintf(buff, "%s/%s_%d.weights", backup_directory, abase, i); save_weights(anet, buff); } if(i%100==0){ char buff[256]; sprintf(buff, "%s/%s.backup", backup_directory, gbase); save_weights(gnet, buff); sprintf(buff, "%s/%s.backup", backup_directory, abase); save_weights(anet, buff); } } #endif }