softmax_layer.c

#include "softmax_layer.h"
#include "blas.h"
#include "dark_cuda.h"
#include "utils.h"
#include "blas.h"

#include <float.h>
#include <math.h>
#include <stdlib.h>
#include <stdio.h>
#include <assert.h>

#define SECRET_NUM -1234

void softmax_tree(float *input, int batch, int inputs, float temp, tree *hierarchy, float *output)
{
	int b;
	for (b = 0; b < batch; ++b) {
		int i;
		int count = 0;
		for (i = 0; i < hierarchy->groups; ++i) {
			int group_size = hierarchy->group_size[i];
			softmax(input + b*inputs + count, group_size, temp, output + b*inputs + count, 1);
			count += group_size;
		}
	}
}

softmax_layer make_softmax_layer(int batch, int inputs, int groups)
{
    assert(inputs%groups == 0);
    fprintf(stderr, "softmax                                        %4d\n",  inputs);
    softmax_layer l = { (LAYER_TYPE)0 };
    l.type = SOFTMAX;
    l.batch = batch;
    l.groups = groups;
    l.inputs = inputs;
    l.outputs = inputs;
    l.loss = (float*)xcalloc(inputs * batch, sizeof(float));
    l.output = (float*)xcalloc(inputs * batch, sizeof(float));
    l.delta = (float*)xcalloc(inputs * batch, sizeof(float));
    l.cost = (float*)xcalloc(1, sizeof(float));

    l.forward = forward_softmax_layer;
    l.backward = backward_softmax_layer;
#ifdef GPU
    l.forward_gpu = forward_softmax_layer_gpu;
    l.backward_gpu = backward_softmax_layer_gpu;

    l.output_gpu = cuda_make_array(l.output, inputs*batch);
    l.loss_gpu = cuda_make_array(l.loss, inputs*batch);
    l.delta_gpu = cuda_make_array(l.delta, inputs*batch);
#endif
    return l;
}

void forward_softmax_layer(const softmax_layer l, network_state net)
{
    if(l.softmax_tree){
        int i;
        int count = 0;
        for (i = 0; i < l.softmax_tree->groups; ++i) {
            int group_size = l.softmax_tree->group_size[i];
            softmax_cpu(net.input + count, group_size, l.batch, l.inputs, 1, 0, 1, l.temperature, l.output + count);
            count += group_size;
        }
    } else {
        softmax_cpu(net.input, l.inputs/l.groups, l.batch, l.inputs, l.groups, l.inputs/l.groups, 1, l.temperature, l.output);
    }

    if(net.truth && !l.noloss){
        softmax_x_ent_cpu(l.batch*l.inputs, l.output, net.truth, l.delta, l.loss);
        l.cost[0] = sum_array(l.loss, l.batch*l.inputs);
    }
}

void backward_softmax_layer(const softmax_layer l, network_state net)
{
    axpy_cpu(l.inputs*l.batch, 1, l.delta, 1, net.delta, 1);
}

#ifdef GPU

void pull_softmax_layer_output(const softmax_layer layer)
{
    cuda_pull_array(layer.output_gpu, layer.output, layer.inputs*layer.batch);
}

void forward_softmax_layer_gpu(const softmax_layer l, network_state net)
{
    if(l.softmax_tree){
		softmax_tree_gpu(net.input, 1, l.batch, l.inputs, l.temperature, l.output_gpu, *l.softmax_tree);
		/*
		int i;
		int count = 0;
		for (i = 0; i < l.softmax_tree->groups; ++i) {
		int group_size = l.softmax_tree->group_size[i];
		softmax_gpu(net.input_gpu + count, group_size, l.batch, l.inputs, 1, 0, 1, l.temperature, l.output_gpu + count);
		count += group_size;
		}
		*/
    } else {
        if(l.spatial){
			softmax_gpu_new_api(net.input, l.c, l.batch*l.c, l.inputs/l.c, l.w*l.h, 1, l.w*l.h, 1, l.output_gpu);
        }else{
			softmax_gpu_new_api(net.input, l.inputs/l.groups, l.batch, l.inputs, l.groups, l.inputs/l.groups, 1, l.temperature, l.output_gpu);
        }
    }
    if(net.truth && !l.noloss){
        softmax_x_ent_gpu(l.batch*l.inputs, l.output_gpu, net.truth, l.delta_gpu, l.loss_gpu);
        if(l.softmax_tree){
			mask_gpu_new_api(l.batch*l.inputs, l.delta_gpu, SECRET_NUM, net.truth, 0);
			mask_gpu_new_api(l.batch*l.inputs, l.loss_gpu, SECRET_NUM, net.truth, 0);
        }
        cuda_pull_array(l.loss_gpu, l.loss, l.batch*l.inputs);
        l.cost[0] = sum_array(l.loss, l.batch*l.inputs);
    }
}

void backward_softmax_layer_gpu(const softmax_layer layer, network_state state)
{
	axpy_ongpu(layer.batch*layer.inputs, state.net.loss_scale, layer.delta_gpu, 1, state.delta, 1);
}

#endif

// -------------------------------------

// Supervised Contrastive Learning: https://arxiv.org/pdf/2004.11362.pdf
contrastive_layer make_contrastive_layer(int batch, int w, int h, int c, int classes, int inputs, layer *yolo_layer)
{
    contrastive_layer l = { (LAYER_TYPE)0 };
    l.type = CONTRASTIVE;
    l.batch = batch;
    l.inputs = inputs;
    l.w = w;
    l.h = h;
    l.c = c;
    l.temperature = 1;

    l.max_boxes = 0;
    if (yolo_layer) {
        l.detection = 1;
        l.max_boxes = yolo_layer->max_boxes;
        l.labels = yolo_layer->labels;  // track id
        l.class_ids = yolo_layer->class_ids;  // class_ids
        l.n = yolo_layer->n;            // num of embeddings per cell = num of anchors
        l.classes = yolo_layer->classes;// num of classes
        classes = l.classes;
        l.embedding_size = l.inputs / (l.n*l.h*l.w);
        l.truths = yolo_layer->truths;
        if (l.embedding_size != yolo_layer->embedding_size) {
            printf(" Error: [contrastive] embedding_size=%d isn't equal to [yolo] embedding_size=%d. They should use the same [convolutional] layer \n", l.embedding_size, yolo_layer->embedding_size);
            getchar();
            exit(0);
        }
        if (l.inputs % (l.n*l.h*l.w) != 0) {
            printf(" Warning: filters= number in the previous (embedding) layer isn't divisable by number of anchors %d \n", l.n);
            getchar();
        }
    }
    else {
        l.detection = 0;
        l.labels = (int*)xcalloc(l.batch, sizeof(int)); // labels
        l.n = 1;                                        // num of embeddings per cell
        l.classes = classes;                            // num of classes
        l.embedding_size = l.c;
    }
    l.outputs = inputs;

    l.loss = (float*)xcalloc(1, sizeof(float));
    l.output = (float*)xcalloc(inputs * batch, sizeof(float));
    l.delta = (float*)xcalloc(inputs * batch, sizeof(float));
    l.cost = (float*)xcalloc(1, sizeof(float));

    const size_t step = l.batch*l.n*l.h*l.w;
    l.cos_sim = NULL;
    l.exp_cos_sim = NULL;
    l.p_constrastive = NULL;
    if (!l.detection) {
        l.cos_sim = (float*)xcalloc(step*step, sizeof(float));
        l.exp_cos_sim = (float*)xcalloc(step*step, sizeof(float));
        l.p_constrastive = (float*)xcalloc(step*step, sizeof(float));
    }
    //l.p_constrastive = (float*)xcalloc(step*step, sizeof(float));
    //l.contrast_p_size = (int*)xcalloc(1, sizeof(int));
    //*l.contrast_p_size = step;
    //l.contrast_p = (contrastive_params*)xcalloc(*l.contrast_p_size, sizeof(contrastive_params));

    l.forward = forward_contrastive_layer;
    l.backward = backward_contrastive_layer;
#ifdef GPU
    l.forward_gpu = forward_contrastive_layer_gpu;
    l.backward_gpu = backward_contrastive_layer_gpu;

    l.output_gpu = cuda_make_array(l.output, inputs*batch);
    l.delta_gpu = cuda_make_array(l.delta, inputs*batch);

    const int max_contr_size = (l.max_boxes*l.batch)*(l.max_boxes*l.batch) * sizeof(contrastive_params)/4;
    printf(" max_contr_size = %d MB \n", max_contr_size / (1024*1024));
    l.contrast_p_gpu = (contrastive_params *)cuda_make_array(NULL, max_contr_size);
#endif
    fprintf(stderr, "contrastive %4d x%4d x%4d x emb_size %4d x batch: %4d  classes = %4d, step = %4d \n", w, h, l.n, l.embedding_size, batch, l.classes, step);
    if(l.detection) fprintf(stderr, "detection \n");
    return l;
}

static inline float clip_value(float val, const float max_val)
{
    if (val > max_val) {
        //printf("\n val = %f > max_val = %f \n", val, max_val);
        val = max_val;
    }
    else if (val < -max_val) {
        //printf("\n val = %f < -max_val = %f \n", val, -max_val);
        val = -max_val;
    }
    return val;
}

void forward_contrastive_layer(contrastive_layer l, network_state state)
{
    if (!state.train) return;
    const float truth_thresh = state.net.label_smooth_eps;

    const int mini_batch = l.batch / l.steps;

    int b, n, w, h;
    fill_cpu(l.batch*l.inputs, 0, l.delta, 1);

    if (!l.detection) {

        for (b = 0; b < l.batch; ++b) {
            if (state.net.adversarial) l.labels[b] = b % 2;
            else l.labels[b] = b / 2;
        }

        // set labels
        for (b = 0; b < l.batch; ++b) {
            for (h = 0; h < l.h; ++h) {
                for (w = 0; w < l.w; ++w)
                {
                    // find truth with max prob (only 1 label even if mosaic is used)
                    float max_truth = 0;
                    int n;
                    for (n = 0; n < l.classes; ++n) {
                        const float truth_prob = state.truth[b*l.classes + n];
                        //printf(" truth_prob = %f, ", truth_prob);
                        //if (truth_prob > max_truth)
                        if (truth_prob > truth_thresh)
                        {
                            //printf(" truth_prob = %f, max_truth = %f, n = %d; ", truth_prob, max_truth, n);
                            max_truth = truth_prob;
                            l.labels[b] = n;
                        }
                    }
                    //printf(", l.labels[b] = %d ", l.labels[b]);
                }
            }
        }

    }
    //printf("\n\n");

    // set pointers to features
    float **z = (float**)xcalloc(l.batch*l.n*l.h*l.w, sizeof(float*));

    for (b = 0; b < l.batch; ++b) {
        for (n = 0; n < l.n; ++n) {
            for (h = 0; h < l.h; ++h) {
                for (w = 0; w < l.w; ++w)
                {
                    const int z_index = b*l.n*l.h*l.w + n*l.h*l.w + h*l.w + w;
                    if (l.labels[z_index] < 0) continue;

                    //const int input_index = b*l.inputs + n*l.embedding_size*l.h*l.w + h*l.w + w;
                    //float *ptr = state.input + input_index;
                    //z[z_index] = ptr;

                    z[z_index] = (float*)xcalloc(l.embedding_size, sizeof(float));
                    get_embedding(state.input, l.w, l.h, l.c, l.embedding_size, w, h, n, b, z[z_index]);
                }
            }
        }
    }

    int b2, n2, h2, w2;
    int contrast_p_index = 0;

    const size_t step = l.batch*l.n*l.h*l.w;
    size_t contrast_p_size = step;
    if (!l.detection) contrast_p_size = l.batch*l.batch;
    contrastive_params *contrast_p = (contrastive_params*)xcalloc(contrast_p_size, sizeof(contrastive_params));

    float *max_sim_same = (float *)xcalloc(l.batch*l.inputs, sizeof(float));
    float *max_sim_diff = (float *)xcalloc(l.batch*l.inputs, sizeof(float));
    fill_cpu(l.batch*l.inputs, -10, max_sim_same, 1);
    fill_cpu(l.batch*l.inputs, -10, max_sim_diff, 1);

    // precalculate cosine similiraty
    for (b = 0; b < l.batch; ++b) {
        for (n = 0; n < l.n; ++n) {
            for (h = 0; h < l.h; ++h) {
                for (w = 0; w < l.w; ++w)
                {
                    const int z_index = b*l.n*l.h*l.w + n*l.h*l.w + h*l.w + w;
                    if (l.labels[z_index] < 0) continue;

                    for (b2 = 0; b2 < l.batch; ++b2) {
                        for (n2 = 0; n2 < l.n; ++n2) {
                            for (h2 = 0; h2 < l.h; ++h2) {
                                for (w2 = 0; w2 < l.w; ++w2)
                                {
                                    const int z_index2 = b2*l.n*l.h*l.w + n2*l.h*l.w + h2*l.w + w2;
                                    if (l.labels[z_index2] < 0) continue;
                                    if (z_index == z_index2) continue;
                                    if (l.detection)
                                        if (l.class_ids[z_index] != l.class_ids[z_index2]) continue;

                                    const int time_step_i = b / mini_batch;
                                    const int time_step_j = b2 / mini_batch;
                                    if (time_step_i != time_step_j) continue;

                                    const size_t step = l.batch*l.n*l.h*l.w;

                                    const float sim = cosine_similarity(z[z_index], z[z_index2], l.embedding_size);
                                    const float exp_sim = expf(sim / l.temperature);
                                    if (!l.detection) {
                                        l.cos_sim[z_index*step + z_index2] = sim;
                                        l.exp_cos_sim[z_index*step + z_index2] = exp_sim;
                                    }

                                    // calc good sim
                                    if (l.labels[z_index] == l.labels[z_index2] && max_sim_same[z_index] < sim) max_sim_same[z_index] = sim;
                                    if (l.labels[z_index] != l.labels[z_index2] && max_sim_diff[z_index] < sim) max_sim_diff[z_index] = sim;
                                    //printf(" z_i = %d, z_i2 = %d, l = %d, l2 = %d, sim = %f \n", z_index, z_index2, l.labels[z_index], l.labels[z_index2], sim);

                                    contrast_p[contrast_p_index].sim = sim;
                                    contrast_p[contrast_p_index].exp_sim = exp_sim;
                                    contrast_p[contrast_p_index].i = z_index;
                                    contrast_p[contrast_p_index].j = z_index2;
                                    contrast_p[contrast_p_index].time_step_i = time_step_i;
                                    contrast_p[contrast_p_index].time_step_j = time_step_j;
                                    contrast_p_index++;
                                    //printf(" contrast_p_index = %d, contrast_p_size = %d \n", contrast_p_index, contrast_p_size);
                                    if ((contrast_p_index+1) >= contrast_p_size) {
                                        contrast_p_size = contrast_p_index + 1;
                                        //printf(" contrast_p_size = %d, z_index = %d, z_index2 = %d \n", contrast_p_size, z_index, z_index2);
                                        contrast_p = (contrastive_params*)xrealloc(contrast_p, contrast_p_size * sizeof(contrastive_params));
                                    }

                                    if (sim > 1.001 || sim < -1.001) {
                                        printf(" sim = %f, ", sim); getchar();
                                    }
                                }
                            }
                        }
                    }
                }
            }
        }
    }

    // calc contrastive accuracy
    int i;
    int good_sims = 0, all_sims = 0, same_sim = 0, diff_sim = 0;
    for (i = 0; i < l.batch*l.inputs; ++i) {
        if (max_sim_same[i] >= -1 && max_sim_diff[i] >= -1) {
            if (max_sim_same[i] >= -1) same_sim++;
            if (max_sim_diff[i] >= -1) diff_sim++;
            ++all_sims;
            //printf(" max_sim_diff[i] = %f, max_sim_same[i] = %f \n", max_sim_diff[i], max_sim_same[i]);
            if (max_sim_diff[i] < max_sim_same[i]) good_sims++;
        }
    }
    if (all_sims > 0) {
        *l.loss = 100 * good_sims / all_sims;
    }
    else *l.loss = -1;
    printf(" Contrast accuracy = %f %%, all = %d, good = %d, same = %d, diff = %d \n", *l.loss, all_sims, good_sims, same_sim, diff_sim);
    free(max_sim_same);
    free(max_sim_diff);


    /*
    // show near sim
    float good_contrast = 0;
    for (b = 0; b < l.batch; b += 2) {
        float same = l.cos_sim[b*l.batch + b];
        float aug = l.cos_sim[b*l.batch + b + 1];
        float diff = l.cos_sim[b*l.batch + b + 2];
        good_contrast += (aug > diff);
        //printf(" l.labels[b] = %d, l.labels[b+1] = %d, l.labels[b+2] = %d, b = %d \n", l.labels[b], l.labels[b + 1], l.labels[b + 2], b);
        //printf(" same = %f, aug = %f, diff = %f, (aug > diff) = %d \n", same, aug, diff, (aug > diff));
    }
    *l.loss = 100 * good_contrast / (l.batch / 2);
    printf(" Contrast accuracy = %f %% \n", *l.loss);
    */

    /*
    // precalculate P_contrastive
    for (b = 0; b < l.batch; ++b) {
        int b2;
        for (b2 = 0; b2 < l.batch; ++b2) {
            if (b != b2) {
                const float P = P_constrastive(b, b2, l.labels, l.batch, z, l.embedding_size, l.temperature, l.cos_sim);
                l.p_constrastive[b*l.batch + b2] = P;
                if (P > 1 || P < -1) {
                    printf(" p = %f, ", P); getchar();
                }
            }
        }
    }
    */


    const size_t contr_size = contrast_p_index;

    if (l.detection) {
#ifdef GPU
        const int max_contr_size = (l.max_boxes*l.batch)*(l.max_boxes*l.batch);
        if (max_contr_size < contr_size) {
            printf(" Error: too large number of bboxes: contr_size = %d > max_contr_size  = %d \n", contr_size, max_contr_size);
            exit(0);
        }
        int *labels = NULL;
        if (contr_size > 2) {
            cuda_push_array((float *)l.contrast_p_gpu, (float *)contrast_p, contr_size * sizeof(contrastive_params) / 4);
            P_constrastive_f_det_gpu(labels, l.embedding_size, l.temperature, l.contrast_p_gpu, contr_size);
            cuda_pull_array((float *)l.contrast_p_gpu, (float *)contrast_p, contr_size * sizeof(contrastive_params) / 4);
        }
#else   // GPU
        int k;
        //#pragma omp parallel for
        for (k = 0; k < contr_size; ++k) {
            contrast_p[k].P = P_constrastive_f_det(k, l.labels, z, l.embedding_size, l.temperature, contrast_p, contr_size);
        }
#endif  // GPU
    }
    else {
        // precalculate P-contrastive
        for (b = 0; b < l.batch; ++b) {
            for (n = 0; n < l.n; ++n) {
                for (h = 0; h < l.h; ++h) {
                    for (w = 0; w < l.w; ++w)
                    {
                        const int z_index = b*l.n*l.h*l.w + n*l.h*l.w + h*l.w + w;
                        if (l.labels[z_index] < 0) continue;

                        for (b2 = 0; b2 < l.batch; ++b2) {
                            for (n2 = 0; n2 < l.n; ++n2) {
                                for (h2 = 0; h2 < l.h; ++h2) {
                                    for (w2 = 0; w2 < l.w; ++w2)
                                    {
                                        const int z_index2 = b2*l.n*l.h*l.w + n2*l.h*l.w + h2*l.w + w2;
                                        if (l.labels[z_index2] < 0) continue;
                                        if (z_index == z_index2) continue;
                                        if (l.detection)
                                            if (l.class_ids[z_index] != l.class_ids[z_index2]) continue;

                                        const int time_step_i = b / mini_batch;
                                        const int time_step_j = b2 / mini_batch;
                                        if (time_step_i != time_step_j) continue;

                                        const size_t step = l.batch*l.n*l.h*l.w;

                                        float P = -10;
                                        if (l.detection) {
                                            P = P_constrastive_f(z_index, z_index2, l.labels, z, l.embedding_size, l.temperature, contrast_p, contr_size);
                                        }
                                        else {
                                            P = P_constrastive(z_index, z_index2, l.labels, step, z, l.embedding_size, l.temperature, l.cos_sim, l.exp_cos_sim);
                                            l.p_constrastive[z_index*step + z_index2] = P;
                                        }

                                        int q;
                                        for (q = 0; q < contr_size; ++q)
                                            if (contrast_p[q].i == z_index && contrast_p[q].j == z_index2) {
                                                contrast_p[q].P = P;
                                                break;
                                            }

                                        //if (q == contr_size) getchar();


                                        //if (P > 1 || P < -1) {
                                        //    printf(" p = %f, z_index = %d, z_index2 = %d ", P, z_index, z_index2); getchar();
                                        //}
                                    }
                                }
                            }
                        }
                    }
                }
            }
        }
    }


    // calc deltas
    #pragma omp parallel for
    for (b = 0; b < l.batch; ++b) {
        for (n = 0; n < l.n; ++n) {
            for (h = 0; h < l.h; ++h) {
                for (w = 0; w < l.w; ++w)
                {
                    const int z_index = b*l.n*l.h*l.w + n*l.h*l.w + h*l.w + w;
                    const size_t step = l.batch*l.n*l.h*l.w;
                    if (l.labels[z_index] < 0) continue;

                    const int delta_index = b*l.embedding_size*l.n*l.h*l.w + n*l.embedding_size*l.h*l.w + h*l.w + w;
                    const int wh = l.w*l.h;

                    if (l.detection) {
                        // detector

                        // positive
                        grad_contrastive_loss_positive_f(z_index, l.class_ids, l.labels, step, z, l.embedding_size, l.temperature, l.delta + delta_index, wh, contrast_p, contr_size);

                        // negative
                        grad_contrastive_loss_negative_f(z_index, l.class_ids, l.labels, step, z, l.embedding_size, l.temperature, l.delta + delta_index, wh, contrast_p, contr_size, l.contrastive_neg_max);
                    }
                    else {
                        // classifier

                        // positive
                        grad_contrastive_loss_positive(z_index, l.labels, step, z, l.embedding_size, l.temperature, l.cos_sim, l.p_constrastive, l.delta + delta_index, wh);

                        // negative
                        grad_contrastive_loss_negative(z_index, l.labels, step, z, l.embedding_size, l.temperature, l.cos_sim, l.p_constrastive, l.delta + delta_index, wh);
                    }

                }
            }
        }
    }

    scal_cpu(l.inputs * l.batch, l.cls_normalizer, l.delta, 1);

    for (i = 0; i < l.inputs * l.batch; ++i) {
        l.delta[i] = clip_value(l.delta[i], l.max_delta);
    }

    *(l.cost) = pow(mag_array(l.delta, l.inputs * l.batch), 2);
    if (state.net.adversarial) {
        printf(" adversarial contrastive loss = %f \n\n", *(l.cost));
    }
    else {
        printf(" contrastive loss = %f \n\n", *(l.cost));
    }

    for (b = 0; b < l.batch; ++b) {
        for (n = 0; n < l.n; ++n) {
            for (h = 0; h < l.h; ++h) {
                for (w = 0; w < l.w; ++w)
                {
                    const int z_index = b*l.n*l.h*l.w + n*l.h*l.w + h*l.w + w;
                    //if (l.labels[z_index] < 0) continue;
                    if (z[z_index]) free(z[z_index]);
                }
            }
        }
    }

    free(contrast_p);
    free(z);
}

void backward_contrastive_layer(contrastive_layer l, network_state state)
{
    axpy_cpu(l.inputs*l.batch, 1, l.delta, 1, state.delta, 1);
}


#ifdef GPU

void pull_contrastive_layer_output(const contrastive_layer l)
{
    cuda_pull_array(l.output_gpu, l.output, l.inputs*l.batch);
}

void push_contrastive_layer_output(const contrastive_layer l)
{
    cuda_push_array(l.delta_gpu, l.delta, l.inputs*l.batch);
}


void forward_contrastive_layer_gpu(contrastive_layer l, network_state state)
{
    simple_copy_ongpu(l.batch*l.inputs, state.input, l.output_gpu);
    if (!state.train) return;

    float *in_cpu = (float *)xcalloc(l.batch*l.inputs, sizeof(float));
    cuda_pull_array(l.output_gpu, l.output, l.batch*l.outputs);
    memcpy(in_cpu, l.output, l.batch*l.outputs * sizeof(float));
    float *truth_cpu = 0;
    if (state.truth) {
        int num_truth = l.batch*l.classes;
        if (l.detection) num_truth = l.batch*l.truths;
        truth_cpu = (float *)xcalloc(num_truth, sizeof(float));
        cuda_pull_array(state.truth, truth_cpu, num_truth);
    }
    network_state cpu_state = state;
    cpu_state.net = state.net;
    cpu_state.index = state.index;
    cpu_state.train = state.train;
    cpu_state.truth = truth_cpu;
    cpu_state.input = in_cpu;

    forward_contrastive_layer(l, cpu_state);
    cuda_push_array(l.delta_gpu, l.delta, l.batch*l.outputs);

    free(in_cpu);
    if (cpu_state.truth) free(cpu_state.truth);
}

void backward_contrastive_layer_gpu(contrastive_layer layer, network_state state)
{
    axpy_ongpu(layer.batch*layer.inputs, state.net.loss_scale, layer.delta_gpu, 1, state.delta, 1);
}

#endif