 }
 
 
+template<typename T>
+__device__ __host__
+void swap(T& a, T& b)
+{
+    T tmp = b;
+    b = a;
+    a = tmp;
+}
+
+
 // sorts an array of size_t and mirrors the same permutation to two arrays of double
 __device__ __host__
 void sort_sparse_expr(expr_t* const expr)

    else
    {
        // if a points to the temporaries created in merge_sort, swap, then memcpy from b to a, then free the temporaries
        std::swap(a, b);
        std::swap(mirror_1_a, mirror_1_b);
        std::swap(mirror_2_a, mirror_2_b);

        memcpy(a, b, length*sizeof(size_t));
        memcpy(mirror_1_a, mirror_1_b, length*sizeof(double));

    create_sparse_expr_device_neuron<<<1, 1>>>(neurons, neuron_index, inf_coeff_tmp, sup_coeff_tmp, dim_tmp, src->inf_cst, src->sup_cst, src->size, src->type);
}

void device_layer_create_sparse_exprs(neuron_t** out_neurons, const double* filter_weights, const double* filter_bias,
                               const size_t* input_size, const size_t* output_size, const size_t* filter_size, const size_t* strides,
                                const bool has_bias, const long int pad_top, const long int pad_left, const size_t num_pixels)
{
    for(size_t out_x = 0; out_x < output_size[0]; out_x++)
    {
        for(size_t out_y = 0; out_y < output_size[1]; out_y++)
        {
            for(size_t out_z = 0; out_z < output_size[2]; out_z++)
            {
                const size_t mat_x = out_x*output_size[1]*output_size[2] + out_y*output_size[2] + out_z;
                const size_t num_coeff = input_size[2]*filter_size[0]*filter_size[1];
