From abf96380a2bcdde795bc88664740dd835ce3f65c Mon Sep 17 00:00:00 2001 From: milakov Date: Sat, 27 Jul 2013 17:29:11 +0400 Subject: [PATCH] Regularization "Upper bound on L2 norm of the incoming weight vector for each output neuron" added --- .../convolution_weight_vector_bound_cuda.cu | 175 ++++++++++++++++++ .../convolution_weight_vector_bound_cuda.h | 54 ++++++ nnforge/cuda/cuda.cpp | 5 + nnforge/cuda/network_updater_cuda.cu | 36 +++- nnforge/cuda/network_updater_cuda.h | 4 + nnforge/cuda/network_updater_cuda_factory.cpp | 9 +- nnforge/cuda/network_updater_cuda_factory.h | 3 +- nnforge/cuda/weight_vector_bound_cuda.cpp | 90 +++++++++ nnforge/cuda/weight_vector_bound_cuda.h | 81 ++++++++ .../cuda/weight_vector_bound_cuda_factory.cpp | 51 +++++ .../cuda/weight_vector_bound_cuda_factory.h | 50 +++++ nnforge/network_updater.cpp | 25 ++- nnforge/network_updater.h | 8 +- nnforge/network_updater_factory.h | 5 +- nnforge/neural_network_toolset.cpp | 15 +- nnforge/neural_network_toolset.h | 3 + .../convolution_weight_vector_bound_plain.cpp | 76 ++++++++ .../convolution_weight_vector_bound_plain.h | 60 ++++++ nnforge/plain/network_updater_plain.cpp | 27 ++- nnforge/plain/network_updater_plain.h | 3 + .../plain/network_updater_plain_factory.cpp | 9 +- nnforge/plain/network_updater_plain_factory.h | 3 +- nnforge/plain/plain.cpp | 8 +- nnforge/plain/weight_vector_bound_plain.cpp | 31 ++++ nnforge/plain/weight_vector_bound_plain.h | 56 ++++++ .../weight_vector_bound_plain_factory.cpp | 48 +++++ .../plain/weight_vector_bound_plain_factory.h | 48 +++++ nnforge/weight_vector_bound.cpp | 29 +++ nnforge/weight_vector_bound.h | 30 +++ 29 files changed, 1022 insertions(+), 20 deletions(-) create mode 100644 nnforge/cuda/convolution_weight_vector_bound_cuda.cu create mode 100644 nnforge/cuda/convolution_weight_vector_bound_cuda.h create mode 100644 nnforge/cuda/weight_vector_bound_cuda.cpp create mode 100644 nnforge/cuda/weight_vector_bound_cuda.h create mode 100644 nnforge/cuda/weight_vector_bound_cuda_factory.cpp create mode 100644 nnforge/cuda/weight_vector_bound_cuda_factory.h create mode 100644 nnforge/plain/convolution_weight_vector_bound_plain.cpp create mode 100644 nnforge/plain/convolution_weight_vector_bound_plain.h create mode 100644 nnforge/plain/weight_vector_bound_plain.cpp create mode 100644 nnforge/plain/weight_vector_bound_plain.h create mode 100644 nnforge/plain/weight_vector_bound_plain_factory.cpp create mode 100644 nnforge/plain/weight_vector_bound_plain_factory.h create mode 100644 nnforge/weight_vector_bound.cpp create mode 100644 nnforge/weight_vector_bound.h diff --git a/nnforge/cuda/convolution_weight_vector_bound_cuda.cu b/nnforge/cuda/convolution_weight_vector_bound_cuda.cu new file mode 100644 index 0000000..cadb1b0 --- /dev/null +++ b/nnforge/cuda/convolution_weight_vector_bound_cuda.cu @@ -0,0 +1,175 @@ +/* + * Copyright 2011-2013 Maxim Milakov + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#include "convolution_weight_vector_bound_cuda.h" + +#include "../convolution_layer.h" + +extern __shared__ float arr[]; +template +__global__ void convolution_normalize_weights_to_max_l2_norm_kernel( + float * __restrict weights, + const float * __restrict weights_read_copy, + float max_l2_norm_squared, + int incoming_weight_count_per_output_neuron, + int output_feature_map_count, + int min_iteration_count) +{ + int thread_id = threadIdx.x; + int output_feature_map_id = blockIdx.y; + int entry_id = blockIdx.z; + int threadblock_size = blockDim.x; + + int base_weight_id = (entry_id * output_feature_map_count + output_feature_map_id) * incoming_weight_count_per_output_neuron; + float current_val; + float sum = 0.0F; + int current_weight_id = thread_id; + for(int i = 0; i < min_iteration_count; ++i) + { + current_val = weights_read_copy[base_weight_id + current_weight_id]; + sum += current_val * current_val; + current_weight_id += threadblock_size; + } + if (current_weight_id < incoming_weight_count_per_output_neuron) + { + current_val = weights_read_copy[base_weight_id + current_weight_id]; + sum += current_val * current_val; + } + arr[thread_id] = sum; + __syncthreads(); + + int t_add_elems = threadblock_size >> 1; + int t_working_elems = (threadblock_size + 1) >> 1; + while (t_add_elems > 0) + { + if (thread_id < t_add_elems) + arr[thread_id] += arr[thread_id + t_working_elems]; + t_add_elems = t_working_elems >> 1; + t_working_elems = (t_working_elems + 1) >> 1; + __syncthreads(); + } + + sum = arr[0]; + if (sum <= max_l2_norm_squared) + return; + + float mult = rsqrtf(__fdividef(sum, max_l2_norm_squared)); + + if (single_item_per_thread) + { + if (thread_id < incoming_weight_count_per_output_neuron) + weights[base_weight_id + thread_id] = current_val * mult; + } + else + { + int current_weight_id = thread_id; + for(int i = 0; i < min_iteration_count; ++i) + { + weights[base_weight_id + current_weight_id] += weights_read_copy[base_weight_id + current_weight_id] * mult; + current_weight_id += threadblock_size; + } + if (current_weight_id < incoming_weight_count_per_output_neuron) + weights[base_weight_id + current_weight_id] += weights_read_copy[base_weight_id + current_weight_id] * mult; + } +} + +namespace nnforge +{ + namespace cuda + { + convolution_weight_vector_bound_cuda::convolution_weight_vector_bound_cuda() + { + } + + convolution_weight_vector_bound_cuda::~convolution_weight_vector_bound_cuda() + { + } + + const boost::uuids::uuid& convolution_weight_vector_bound_cuda::get_uuid() const + { + return convolution_layer::layer_guid; + } + + void convolution_weight_vector_bound_cuda::enqueue_normalize_weights( + cudaStream_t stream_id, + const weight_vector_bound& bound, + const std::vector& data, + const std::vector& additional_buffers, + unsigned int entry_count) + { + int threadblock_size = get_threadblock_size(incoming_weight_count_per_output_neuron); + dim3 grid_size(1, output_feature_map_count, entry_count); + dim3 block_size(threadblock_size, 1, 1); + int min_iteration_count = incoming_weight_count_per_output_neuron / threadblock_size; + int smem_size = threadblock_size * sizeof(float); + float max_l2_norm_squared = bound.max_l2_norm * bound.max_l2_norm; + + if (incoming_weight_count_per_output_neuron <= threadblock_size) + { + convolution_normalize_weights_to_max_l2_norm_kernel<<>>( + *data[0], + *data[0], + max_l2_norm_squared, + incoming_weight_count_per_output_neuron, + output_feature_map_count, + min_iteration_count); + } + else + { + convolution_normalize_weights_to_max_l2_norm_kernel<<>>( + *data[0], + *data[0], + max_l2_norm_squared, + incoming_weight_count_per_output_neuron, + output_feature_map_count, + min_iteration_count); + } + } + + std::tr1::shared_ptr convolution_weight_vector_bound_cuda::create_specific() const + { + return std::tr1::shared_ptr(new convolution_weight_vector_bound_cuda()); + } + + void convolution_weight_vector_bound_cuda::weight_vector_bound_configured() + { + std::tr1::shared_ptr layer_derived = std::tr1::dynamic_pointer_cast(layer_schema); + + incoming_weight_count_per_output_neuron = layer_derived->input_feature_map_count; + for(std::vector::const_iterator it = layer_derived->window_sizes.begin(); it != layer_derived->window_sizes.end(); ++it) + incoming_weight_count_per_output_neuron *= *it; + output_feature_map_count = layer_derived->output_feature_map_count; + } + + int convolution_weight_vector_bound_cuda::get_threadblock_size(int incoming_weight_count_per_output_neuron) + { + int threadblock_size; + + if (incoming_weight_count_per_output_neuron < 256) + { + threadblock_size = (incoming_weight_count_per_output_neuron + 32 - 1) / 32 * 32; + } + else + { + int threadblock_count = (incoming_weight_count_per_output_neuron + 256 - 1) / 256; + threadblock_size = (incoming_weight_count_per_output_neuron + threadblock_count - 1) / threadblock_count; + threadblock_size = (threadblock_size + 32 - 1) / 32 * 32; + } + + return threadblock_size; + } + } +} diff --git a/nnforge/cuda/convolution_weight_vector_bound_cuda.h b/nnforge/cuda/convolution_weight_vector_bound_cuda.h new file mode 100644 index 0000000..7c9793a --- /dev/null +++ b/nnforge/cuda/convolution_weight_vector_bound_cuda.h @@ -0,0 +1,54 @@ +/* + * Copyright 2011-2013 Maxim Milakov + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#pragma once + +#include "weight_vector_bound_cuda.h" + +namespace nnforge +{ + namespace cuda + { + class convolution_weight_vector_bound_cuda: public weight_vector_bound_cuda + { + public: + convolution_weight_vector_bound_cuda(); + + virtual ~convolution_weight_vector_bound_cuda(); + + virtual const boost::uuids::uuid& get_uuid() const; + + virtual void enqueue_normalize_weights( + cudaStream_t stream_id, + const weight_vector_bound& bound, + const std::vector& data, + const std::vector& additional_buffers, + unsigned int entry_count); + + protected: + virtual std::tr1::shared_ptr create_specific() const; + + // The method is called when configuration is finished + virtual void weight_vector_bound_configured(); + + int incoming_weight_count_per_output_neuron; + int output_feature_map_count; + + private: + static int get_threadblock_size(int output_neuron_count); + }; + } +} diff --git a/nnforge/cuda/cuda.cpp b/nnforge/cuda/cuda.cpp index a4e7b27..daec2ed 100644 --- a/nnforge/cuda/cuda.cpp +++ b/nnforge/cuda/cuda.cpp @@ -48,6 +48,9 @@ #include "soft_rectified_linear_layer_updater_schema.h" #include "softmax_layer_updater_schema.h" +#include "weight_vector_bound_cuda_factory.h" +#include "convolution_weight_vector_bound_cuda.h" + namespace nnforge { namespace cuda @@ -83,6 +86,8 @@ namespace nnforge single_layer_updater_schema_factory::get_mutable_instance().register_layer_updater_schema(layer_updater_schema_smart_ptr(new rectified_linear_layer_updater_schema())); single_layer_updater_schema_factory::get_mutable_instance().register_layer_updater_schema(layer_updater_schema_smart_ptr(new soft_rectified_linear_layer_updater_schema())); single_layer_updater_schema_factory::get_mutable_instance().register_layer_updater_schema(layer_updater_schema_smart_ptr(new softmax_layer_updater_schema())); + + single_weight_vector_bound_factory::get_mutable_instance().register_weight_vector_bound(weight_vector_bound_cuda_smart_ptr(new convolution_weight_vector_bound_cuda())); } } } diff --git a/nnforge/cuda/network_updater_cuda.cu b/nnforge/cuda/network_updater_cuda.cu index 6dcf5b4..f08f1a8 100644 --- a/nnforge/cuda/network_updater_cuda.cu +++ b/nnforge/cuda/network_updater_cuda.cu @@ -23,6 +23,7 @@ #include "util_cuda.h" #include "cuda_event.h" #include "layer_updater_schema_factory.h" +#include "weight_vector_bound_cuda_factory.h" #include #include @@ -76,8 +77,9 @@ namespace nnforge network_updater_cuda::network_updater_cuda( network_schema_smart_ptr schema, const std::map& layer_to_dropout_rate_map, + const std::map& layer_to_weight_vector_bound_map, cuda_running_configuration_const_smart_ptr cuda_config) - : network_updater(schema, layer_to_dropout_rate_map) + : network_updater(schema, layer_to_dropout_rate_map, layer_to_weight_vector_bound_map) , cuda_config(cuda_config) { const const_layer_list& layer_list = *schema; @@ -100,6 +102,15 @@ namespace nnforge for(const_layer_list::const_iterator it = start_layer_nonempty_weights_iterator; it != layer_list.end(); ++it) updater_schemas.push_back(single_layer_updater_schema_factory::get_const_instance().create_updater_schema_layer(*it, cuda_config)); + for(std::map::const_iterator it = this->layer_to_weight_vector_bound_map.begin(); it != this->layer_to_weight_vector_bound_map.end(); ++it) + { + unsigned int layer_id = it->first; + if (layer_id < testing_layer_count) + throw neural_network_exception((boost::format("Weight vector bound is specified fo layer %1% while it is in testing part (consisting of %2% layers) of the updater") % layer_id % testing_layer_count).str()); + + weight_vector_bounds.insert(std::make_pair(layer_id, single_weight_vector_bound_factory::get_const_instance().create_weight_vector_bound(layer_list[layer_id], cuda_config))); + } + setup_network_cuda(); for(const_layer_testing_schema_list::const_iterator it = testing_schemas.begin(); it != testing_schemas.end(); ++it) @@ -229,6 +240,10 @@ namespace nnforge output_errors = all_buffers.input_errors_buffer; } + std::map > weight_vector_bound_buffers; + for(std::map::const_iterator it = weight_vector_bounds.begin(); it != weight_vector_bounds.end(); ++it) + weight_vector_bound_buffers.insert(std::make_pair(it->first, it->second->allocate_additional_buffers(max_entry_count))); + cuda_linear_buffer_host_smart_ptr input_host_buf(new cuda_linear_buffer_host(input_neuron_count * max_entry_count * input_neuron_elem_size)); unsigned char * input = *input_host_buf; cuda_linear_buffer_host_smart_ptr output_host_buf(new cuda_linear_buffer_host(output_neuron_count * max_entry_count * sizeof(float))); @@ -411,6 +426,19 @@ namespace nnforge input_and_all_buffers_pack_it->second.additional_buffers, input_and_all_buffers_pack_it->second.dynamic_memobjects, updater_entry_count); + + weight_vector_bound_map::iterator bound_it = weight_vector_bounds.find(reverse_layer_id); + if (bound_it != weight_vector_bounds.end()) + { + const weight_vector_bound& bound = layer_to_weight_vector_bound_map.find(reverse_layer_id)->second; + const std::vector& additional_buffers = weight_vector_bound_buffers.find(reverse_layer_id)->second; + bound_it->second->enqueue_normalize_weights( + *command_stream, + bound, + *net_data_it, + additional_buffers, + updater_entry_count); + } } } @@ -628,6 +656,9 @@ namespace nnforge for(std::vector::const_iterator it = updater_list.begin(); it != updater_list.end(); ++it) (*it)->update_buffer_configuration(buffer_configuration, updater_entry_count); + + for(std::map::const_iterator it = weight_vector_bounds.begin(); it != weight_vector_bounds.end(); ++it) + it->second->update_buffer_configuration(buffer_configuration, updater_entry_count); } unsigned int network_updater_cuda::get_max_batch_size() const @@ -637,6 +668,9 @@ namespace nnforge for(std::vector::const_iterator it = updater_list.begin(); it != updater_list.end(); ++it) (*it)->update_buffer_configuration(buffer_configuration); + for(std::map::const_iterator it = weight_vector_bounds.begin(); it != weight_vector_bounds.end(); ++it) + it->second->update_buffer_configuration(buffer_configuration); + return cuda_config->get_max_entry_count(buffer_configuration, 0.5F); } } diff --git a/nnforge/cuda/network_updater_cuda.h b/nnforge/cuda/network_updater_cuda.h index c21ecda..2e039eb 100644 --- a/nnforge/cuda/network_updater_cuda.h +++ b/nnforge/cuda/network_updater_cuda.h @@ -22,6 +22,7 @@ #include "cuda_stream.h" #include "layer_testing_schema.h" #include "layer_updater_schema.h" +#include "weight_vector_bound_cuda.h" namespace nnforge { @@ -33,6 +34,7 @@ namespace nnforge network_updater_cuda( network_schema_smart_ptr schema, const std::map& layer_to_dropout_rate_map, + const std::map& layer_to_weight_vector_bound_map, cuda_running_configuration_const_smart_ptr cuda_config); virtual ~network_updater_cuda(); @@ -89,6 +91,8 @@ namespace nnforge std::vector > updater_schema_data; std::vector updater_list; + weight_vector_bound_map weight_vector_bounds; + static unsigned int max_entry_count_in_single_batch; }; } diff --git a/nnforge/cuda/network_updater_cuda_factory.cpp b/nnforge/cuda/network_updater_cuda_factory.cpp index 6a9d08b..bd134b3 100644 --- a/nnforge/cuda/network_updater_cuda_factory.cpp +++ b/nnforge/cuda/network_updater_cuda_factory.cpp @@ -33,9 +33,14 @@ namespace nnforge network_updater_smart_ptr network_updater_cuda_factory::create( network_schema_smart_ptr schema, - const std::map& layer_to_dropout_rate_map) const + const std::map& layer_to_dropout_rate_map, + const std::map& layer_to_weight_vector_bound_map) const { - return network_updater_smart_ptr(new network_updater_cuda(schema, layer_to_dropout_rate_map, cuda_config)); + return network_updater_smart_ptr(new network_updater_cuda( + schema, + layer_to_dropout_rate_map, + layer_to_weight_vector_bound_map, + cuda_config)); } } } diff --git a/nnforge/cuda/network_updater_cuda_factory.h b/nnforge/cuda/network_updater_cuda_factory.h index 33163a4..f931c8f 100644 --- a/nnforge/cuda/network_updater_cuda_factory.h +++ b/nnforge/cuda/network_updater_cuda_factory.h @@ -32,7 +32,8 @@ namespace nnforge virtual network_updater_smart_ptr create( network_schema_smart_ptr schema, - const std::map& layer_to_dropout_rate_map) const; + const std::map& layer_to_dropout_rate_map, + const std::map& layer_to_weight_vector_bound_map) const; protected: cuda_running_configuration_const_smart_ptr cuda_config; diff --git a/nnforge/cuda/weight_vector_bound_cuda.cpp b/nnforge/cuda/weight_vector_bound_cuda.cpp new file mode 100644 index 0000000..26d1e28 --- /dev/null +++ b/nnforge/cuda/weight_vector_bound_cuda.cpp @@ -0,0 +1,90 @@ +/* + * Copyright 2011-2013 Maxim Milakov + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#include "weight_vector_bound_cuda.h" + +namespace nnforge +{ + namespace cuda + { + weight_vector_bound_cuda::weight_vector_bound_cuda() + { + } + + weight_vector_bound_cuda::~weight_vector_bound_cuda() + { + } + + std::tr1::shared_ptr weight_vector_bound_cuda::create( + const_layer_smart_ptr layer_schema, + cuda_running_configuration_const_smart_ptr cuda_config) const + { + std::tr1::shared_ptr res = create_specific(); + + res->layer_schema = layer_schema; + res->cuda_config = cuda_config; + + res->weight_vector_bound_configured(); + + return res; + } + + void weight_vector_bound_cuda::weight_vector_bound_configured() + { + } + + void weight_vector_bound_cuda::update_buffer_configuration(buffer_cuda_size_configuration& buffer_configuration) const + { + std::vector per_entry_sizes = get_sizes_of_additional_buffers_per_entry(); + for(std::vector::const_iterator it = per_entry_sizes.begin(); it != per_entry_sizes.end(); ++it) + buffer_configuration.add_per_entry_buffer(*it); + + std::vector tex_per_entry = get_linear_addressing_through_texture_per_entry(); + for(std::vector::const_iterator it = tex_per_entry.begin(); it != tex_per_entry.end(); ++it) + buffer_configuration.add_per_entry_linear_addressing_through_texture(*it); + } + + void weight_vector_bound_cuda::update_buffer_configuration( + buffer_cuda_size_configuration& buffer_configuration, + unsigned int updater_entry_count) const + { + std::vector per_entry_sizes = get_sizes_of_additional_buffers_per_entry(); + for(std::vector::const_iterator it = per_entry_sizes.begin(); it != per_entry_sizes.end(); ++it) + buffer_configuration.add_constant_buffer(*it * updater_entry_count); + } + + std::vector weight_vector_bound_cuda::allocate_additional_buffers(unsigned int max_entry_count) + { + std::vector res; + + std::vector per_entry_sizes = get_sizes_of_additional_buffers_per_entry(); + for(std::vector::const_iterator it = per_entry_sizes.begin(); it != per_entry_sizes.end(); ++it) + res.push_back(cuda_linear_buffer_device_smart_ptr(new cuda_linear_buffer_device(*it * max_entry_count))); + + return res; + } + + std::vector weight_vector_bound_cuda::get_sizes_of_additional_buffers_per_entry() const + { + return std::vector(); + } + + std::vector weight_vector_bound_cuda::get_linear_addressing_through_texture_per_entry() const + { + return std::vector(); + } + } +} diff --git a/nnforge/cuda/weight_vector_bound_cuda.h b/nnforge/cuda/weight_vector_bound_cuda.h new file mode 100644 index 0000000..8fa888c --- /dev/null +++ b/nnforge/cuda/weight_vector_bound_cuda.h @@ -0,0 +1,81 @@ +/* + * Copyright 2011-2013 Maxim Milakov + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#pragma once + +#include "../layer.h" +#include "cuda_running_configuration.h" +#include "buffer_cuda_size_configuration.h" +#include "cuda_linear_buffer_device.h" + +#include "../weight_vector_bound.h" + +#include + +namespace nnforge +{ + namespace cuda + { + class weight_vector_bound_cuda + { + public: + virtual ~weight_vector_bound_cuda(); + + std::tr1::shared_ptr create( + const_layer_smart_ptr layer_schema, + cuda_running_configuration_const_smart_ptr cuda_config) const; + + virtual const boost::uuids::uuid& get_uuid() const = 0; + + std::vector allocate_additional_buffers(unsigned int max_entry_count); + + void update_buffer_configuration(buffer_cuda_size_configuration& buffer_configuration) const; + + void update_buffer_configuration( + buffer_cuda_size_configuration& buffer_configuration, + unsigned int updater_entry_count) const; + + virtual void enqueue_normalize_weights( + cudaStream_t stream_id, + const weight_vector_bound& bound, + const std::vector& data, + const std::vector& additional_buffers, + unsigned int entry_count) = 0; + + protected: + weight_vector_bound_cuda(); + + virtual std::tr1::shared_ptr create_specific() const = 0; + + // The method is called when configuration is finished + virtual void weight_vector_bound_configured(); + + virtual std::vector get_sizes_of_additional_buffers_per_entry() const; + + virtual std::vector get_linear_addressing_through_texture_per_entry() const; + + const_layer_smart_ptr layer_schema; + cuda_running_configuration_const_smart_ptr cuda_config; + + private: + weight_vector_bound_cuda(const weight_vector_bound_cuda&); + weight_vector_bound_cuda& operator =(const weight_vector_bound_cuda&); + }; + + typedef std::tr1::shared_ptr weight_vector_bound_cuda_smart_ptr; + typedef std::map weight_vector_bound_map; + } +} diff --git a/nnforge/cuda/weight_vector_bound_cuda_factory.cpp b/nnforge/cuda/weight_vector_bound_cuda_factory.cpp new file mode 100644 index 0000000..50957c9 --- /dev/null +++ b/nnforge/cuda/weight_vector_bound_cuda_factory.cpp @@ -0,0 +1,51 @@ +/* + * Copyright 2011-2013 Maxim Milakov + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#include "weight_vector_bound_cuda_factory.h" +#include "../neural_network_exception.h" + +#include +#include + +namespace nnforge +{ + namespace cuda + { + bool weight_vector_bound_cuda_factory::register_weight_vector_bound(weight_vector_bound_cuda_smart_ptr sample_weight_vector_bound) + { + return sample_weight_vector_bound_map.insert(sample_map::value_type(sample_weight_vector_bound->get_uuid(), sample_weight_vector_bound)).second; + } + + bool weight_vector_bound_cuda_factory::unregister_weight_vector_bound(const boost::uuids::uuid& layer_guid) + { + return sample_weight_vector_bound_map.erase(layer_guid) == 1; + } + + weight_vector_bound_cuda_smart_ptr weight_vector_bound_cuda_factory::create_weight_vector_bound( + const_layer_smart_ptr layer, + cuda_running_configuration_const_smart_ptr cuda_config) const + { + sample_map::const_iterator i = sample_weight_vector_bound_map.find(layer->get_uuid()); + + if (i == sample_weight_vector_bound_map.end()) + throw neural_network_exception((boost::format("No CUDA weight vector bound is registered with id %1%") % layer->get_uuid()).str()); + + return i->second->create( + layer, + cuda_config); + } + } +} diff --git a/nnforge/cuda/weight_vector_bound_cuda_factory.h b/nnforge/cuda/weight_vector_bound_cuda_factory.h new file mode 100644 index 0000000..0c76905 --- /dev/null +++ b/nnforge/cuda/weight_vector_bound_cuda_factory.h @@ -0,0 +1,50 @@ +/* + * Copyright 2011-2013 Maxim Milakov + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#pragma once + +#include "../layer.h" +#include "weight_vector_bound_cuda.h" +#include "cuda_running_configuration.h" + +#include +#include +#include +#include + +namespace nnforge +{ + namespace cuda + { + class weight_vector_bound_cuda_factory + { + public: + bool register_weight_vector_bound(weight_vector_bound_cuda_smart_ptr sample_weight_vector_bound); + + bool unregister_weight_vector_bound(const boost::uuids::uuid& layer_guid); + + weight_vector_bound_cuda_smart_ptr create_weight_vector_bound( + const_layer_smart_ptr layer, + cuda_running_configuration_const_smart_ptr cuda_config) const; + + private: + typedef std::map sample_map; + sample_map sample_weight_vector_bound_map; + }; + + typedef boost::serialization::singleton single_weight_vector_bound_factory; + } +} diff --git a/nnforge/network_updater.cpp b/nnforge/network_updater.cpp index a65c556..66eee4b 100644 --- a/nnforge/network_updater.cpp +++ b/nnforge/network_updater.cpp @@ -27,22 +27,24 @@ namespace nnforge network_updater::network_updater( network_schema_smart_ptr schema, - const std::map& layer_to_dropout_rate_map) + const std::map& layer_to_dropout_rate_map, + const std::map& layer_to_weight_vector_bound_map) : schema(schema) , layer_to_dropout_rate_map(layer_to_dropout_rate_map) , random_uniform_list(1 << random_list_bits) + , layer_to_weight_vector_bound_map(layer_to_weight_vector_bound_map) , profile_mode(false) , gen(rnd::get_random_generator()) { const const_layer_list& layer_list = *schema; unsigned int layer_count = static_cast(layer_list.size()); - for(std::map::const_iterator it = layer_to_dropout_rate_map.begin(); it != layer_to_dropout_rate_map.end(); ++it) - if (it->first >= layer_count) - throw neural_network_exception("Dropout is specified for the layer which doesn't exist in the schema"); std::vector::const_iterator layer_it = schema->get_layers().begin(); for(std::map::const_iterator it = layer_to_dropout_rate_map.begin(); it != layer_to_dropout_rate_map.end(); ++it, ++layer_it) { + if (it->first >= layer_count) + throw neural_network_exception("Dropout is specified for the layer which doesn't exist in the schema"); + float dropout_rate = it->second; if ((dropout_rate < 0.0F) || (dropout_rate >= 1.0F)) @@ -54,6 +56,21 @@ namespace nnforge layer_id_to_dropout_config_map.insert(std::make_pair(it->first, mult)); } } + + for(std::map::iterator it = this->layer_to_weight_vector_bound_map.begin(); it != this->layer_to_weight_vector_bound_map.end(); ++it) + { + if (it->first >= layer_count) + throw neural_network_exception("Weight bound is specified for the layer which doesn't exist in the schema"); + + std::map::const_iterator it2 = layer_id_to_dropout_config_map.find(it->first); + if (it2 != layer_id_to_dropout_config_map.end()) + { + float val = 1.0F; + for(std::map::const_iterator it3 = it2->second.weight_part_to_dropout_direct_multiplier_map.begin(); it3 != it2->second.weight_part_to_dropout_direct_multiplier_map.end(); ++it3) + val = std::min(val, it3->second); + it->second.max_l2_norm = it->second.max_l2_norm / val; + } + } } network_updater::~network_updater() diff --git a/nnforge/network_updater.h b/nnforge/network_updater.h index 0804a9d..cb9a921 100644 --- a/nnforge/network_updater.h +++ b/nnforge/network_updater.h @@ -22,6 +22,7 @@ #include "supervised_data_reader.h" #include "testing_result.h" #include "dropout_layer_config.h" +#include "weight_vector_bound.h" #include #include @@ -54,7 +55,8 @@ namespace nnforge protected: network_updater( network_schema_smart_ptr schema, - const std::map& layer_to_dropout_rate_map); + const std::map& layer_to_dropout_rate_map, + const std::map& layer_to_weight_vector_bound_map); // schema, data and reader are guaranteed to be compatible virtual std::vector actual_update( @@ -73,8 +75,8 @@ namespace nnforge std::map layer_to_dropout_rate_map; layer_configuration_specific_list layer_config_list; std::vector random_uniform_list; - std::map layer_id_to_dropout_config_map; float flops; + std::map layer_to_weight_vector_bound_map; private: network_updater(); @@ -83,6 +85,8 @@ namespace nnforge random_generator gen; static const unsigned int random_list_bits; + + std::map layer_id_to_dropout_config_map; }; typedef std::tr1::shared_ptr network_updater_smart_ptr; diff --git a/nnforge/network_updater_factory.h b/nnforge/network_updater_factory.h index 192328f..6eb78ab 100644 --- a/nnforge/network_updater_factory.h +++ b/nnforge/network_updater_factory.h @@ -17,6 +17,8 @@ #pragma once #include "network_updater.h" +#include "weight_vector_bound.h" + #include namespace nnforge @@ -28,7 +30,8 @@ namespace nnforge virtual network_updater_smart_ptr create( network_schema_smart_ptr schema, - const std::map& layer_to_dropout_rate_map) const = 0; + const std::map& layer_to_dropout_rate_map, + const std::map& layer_to_weight_vector_bound_map) const = 0; protected: network_updater_factory(); diff --git a/nnforge/neural_network_toolset.cpp b/nnforge/neural_network_toolset.cpp index 57faace..b2a7e7c 100644 --- a/nnforge/neural_network_toolset.cpp +++ b/nnforge/neural_network_toolset.cpp @@ -785,7 +785,10 @@ namespace nnforge hessian_calculator_smart_ptr hessian = hessian_factory->create(schema); - network_updater_smart_ptr updater = updater_factory->create(schema, get_dropout_rate_map()); + network_updater_smart_ptr updater = updater_factory->create( + schema, + get_dropout_rate_map(), + get_weight_vector_bound_map()); network_trainer_sdlm trainer( schema, @@ -853,7 +856,10 @@ namespace nnforge schema->read(in); } - network_updater_smart_ptr updater = updater_factory->create(schema, get_dropout_rate_map()); + network_updater_smart_ptr updater = updater_factory->create( + schema, + get_dropout_rate_map(), + get_weight_vector_bound_map()); supervised_data_reader_smart_ptr training_data_reader = get_data_reader_for_training(); @@ -964,6 +970,11 @@ namespace nnforge return std::map(); } + std::map neural_network_toolset::get_weight_vector_bound_map() const + { + return std::map(); + } + std::vector neural_network_toolset::get_data_transformer_list_for_training() const { return std::vector(); diff --git a/nnforge/neural_network_toolset.h b/nnforge/neural_network_toolset.h index a70cfaf..eeb888c 100644 --- a/nnforge/neural_network_toolset.h +++ b/nnforge/neural_network_toolset.h @@ -27,6 +27,7 @@ #include "layer_data_configuration.h" #include "data_transformer.h" #include "data_transformer_util.h" +#include "weight_vector_bound.h" #include @@ -61,6 +62,8 @@ namespace nnforge virtual std::map get_dropout_rate_map() const; + virtual std::map get_weight_vector_bound_map() const; + virtual boost::filesystem::path get_input_data_folder() const; virtual boost::filesystem::path get_working_data_folder() const; diff --git a/nnforge/plain/convolution_weight_vector_bound_plain.cpp b/nnforge/plain/convolution_weight_vector_bound_plain.cpp new file mode 100644 index 0000000..8ae9a2a --- /dev/null +++ b/nnforge/plain/convolution_weight_vector_bound_plain.cpp @@ -0,0 +1,76 @@ +/* + * Copyright 2011-2013 Maxim Milakov + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#include "convolution_weight_vector_bound_plain.h" + +#include "../convolution_layer.h" + +#include + +namespace nnforge +{ + namespace plain + { + convolution_weight_vector_bound_plain::convolution_weight_vector_bound_plain() + { + } + + convolution_weight_vector_bound_plain::~convolution_weight_vector_bound_plain() + { + } + + const boost::uuids::uuid& convolution_weight_vector_bound_plain::get_uuid() const + { + return convolution_layer::layer_guid; + } + + void convolution_weight_vector_bound_plain::normalize_weights( + const weight_vector_bound& bound, + layer_data_list& data, + plain_running_configuration_const_smart_ptr plain_config, + const_layer_smart_ptr layer_schema, + unsigned int updater_count) const + { + std::tr1::shared_ptr layer_derived = std::tr1::dynamic_pointer_cast(layer_schema); + + unsigned int weight_count = 1; + for(std::vector::const_iterator it = layer_derived->window_sizes.begin(); it != layer_derived->window_sizes.end(); ++it) + weight_count *= *it; + const unsigned int weight_block_size = weight_count * layer_derived->input_feature_map_count; + const unsigned int output_feature_map_count = layer_derived->output_feature_map_count; + const int total_workload = output_feature_map_count * updater_count; + const layer_data_list::iterator data_list_it = data.begin(); + const float max_l2_norm_squared = bound.max_l2_norm * bound.max_l2_norm; + const accum_helper_struct accum_helper; + + #pragma omp parallel for default(none) num_threads(plain_config->openmp_thread_count) schedule(guided) + for(int workload_id = 0; workload_id < total_workload; ++workload_id) + { + int entry_id = workload_id / output_feature_map_count; + int output_feature_map_id = workload_id - (entry_id * output_feature_map_count); + + std::vector::iterator weights = (**(data_list_it + entry_id))[0].begin() + (output_feature_map_id * weight_block_size); + + float l2_norm_squared = std::accumulate(weights, weights + weight_block_size, 0.0F, accum_helper); + if (l2_norm_squared > max_l2_norm_squared) + { + float mult = sqrtf(max_l2_norm_squared / l2_norm_squared); + std::transform(weights, weights + weight_block_size, weights, scale_helper_struct(mult)); + } + } + } + } +} diff --git a/nnforge/plain/convolution_weight_vector_bound_plain.h b/nnforge/plain/convolution_weight_vector_bound_plain.h new file mode 100644 index 0000000..029340a --- /dev/null +++ b/nnforge/plain/convolution_weight_vector_bound_plain.h @@ -0,0 +1,60 @@ +/* + * Copyright 2011-2013 Maxim Milakov + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#pragma once + +#include "weight_vector_bound_plain.h" + +namespace nnforge +{ + namespace plain + { + class convolution_weight_vector_bound_plain : public weight_vector_bound_plain + { + public: + convolution_weight_vector_bound_plain(); + + virtual ~convolution_weight_vector_bound_plain(); + + virtual const boost::uuids::uuid& get_uuid() const; + + virtual void normalize_weights( + const weight_vector_bound& bound, + layer_data_list& data, + plain_running_configuration_const_smart_ptr plain_config, + const_layer_smart_ptr layer_schema, + unsigned int updater_count) const; + + private: + struct accum_helper_struct + { + float operator()(float x, float y) {return y * y + x;} + }; + + struct scale_helper_struct + { + scale_helper_struct(float mult) + : mult(mult) + { + } + + float operator()(float x) {return x * mult;} + + float mult; + }; + }; + } +} diff --git a/nnforge/plain/network_updater_plain.cpp b/nnforge/plain/network_updater_plain.cpp index 7eeaa25..ca14cca 100644 --- a/nnforge/plain/network_updater_plain.cpp +++ b/nnforge/plain/network_updater_plain.cpp @@ -22,19 +22,21 @@ #include "layer_tester_plain_factory.h" #include "layer_updater_plain_factory.h" +#include "weight_vector_bound_plain_factory.h" #include "../neural_network_exception.h" namespace nnforge { namespace plain { - unsigned int network_updater_plain::max_entry_count_in_single_batch = 1024; + unsigned int network_updater_plain::max_entry_count_in_single_batch = 5;//1024; network_updater_plain::network_updater_plain( network_schema_smart_ptr schema, const std::map& layer_to_dropout_rate_map, + const std::map& layer_to_weight_vector_bound_map, plain_running_configuration_const_smart_ptr plain_config) - : network_updater(schema, layer_to_dropout_rate_map) + : network_updater(schema, layer_to_dropout_rate_map, layer_to_weight_vector_bound_map) , plain_config(plain_config) { const const_layer_list& layer_list = *schema; @@ -56,6 +58,15 @@ namespace nnforge for(const_layer_list::const_iterator it = start_layer_nonempty_weights_iterator; it != layer_list.end(); ++it) updater_list.push_back(single_layer_updater_plain_factory::get_const_instance().get_updater_plain_layer((*it)->get_uuid())); + + for(std::map::const_iterator it = this->layer_to_weight_vector_bound_map.begin(); it != this->layer_to_weight_vector_bound_map.end(); ++it) + { + unsigned int layer_id = it->first; + if (layer_id < testing_layer_count) + throw neural_network_exception((boost::format("Weight vector bound is specified fo layer %1% while it is in testing part (consisting of %2% layers) of the updater") % layer_id % testing_layer_count).str()); + + weight_vector_bounds.insert(std::make_pair(layer_id, single_weight_vector_bound_factory::get_const_instance().get_updater_plain_layer(layer_list[layer_id]->get_uuid()))); + } } network_updater_plain::~network_updater_plain() @@ -355,6 +366,18 @@ namespace nnforge updater_entry_count, (it == updater_list.rend() - 1) ? input_entry_id : -1); + weight_vector_bound_map::iterator bound_it = weight_vector_bounds.find(reverse_layer_id); + if (bound_it != weight_vector_bounds.end()) + { + const weight_vector_bound& bound = layer_to_weight_vector_bound_map.find(reverse_layer_id)->second; + bound_it->second->normalize_weights( + bound, + *data_it, + plain_config, + *layer_it, + updater_entry_count); + } + output_errors = updater_buffers_it->second.input_errors_buffer; } } diff --git a/nnforge/plain/network_updater_plain.h b/nnforge/plain/network_updater_plain.h index 8ba074b..c4d1fe6 100644 --- a/nnforge/plain/network_updater_plain.h +++ b/nnforge/plain/network_updater_plain.h @@ -22,6 +22,7 @@ #include "plain_running_configuration.h" #include "buffer_plain_size_configuration.h" #include "layer_tester_plain.h" +#include "weight_vector_bound_plain.h" namespace nnforge { @@ -33,6 +34,7 @@ namespace nnforge network_updater_plain( network_schema_smart_ptr schema, const std::map& layer_to_dropout_rate_map, + const std::map& layer_to_weight_vector_bound_map, plain_running_configuration_const_smart_ptr plain_config); ~network_updater_plain(); @@ -65,6 +67,7 @@ namespace nnforge const_layer_tester_plain_list tester_list; const_layer_updater_plain_list updater_list; + weight_vector_bound_map weight_vector_bounds; static unsigned int max_entry_count_in_single_batch; }; diff --git a/nnforge/plain/network_updater_plain_factory.cpp b/nnforge/plain/network_updater_plain_factory.cpp index e0522e9..2fe15a9 100644 --- a/nnforge/plain/network_updater_plain_factory.cpp +++ b/nnforge/plain/network_updater_plain_factory.cpp @@ -33,9 +33,14 @@ namespace nnforge network_updater_smart_ptr network_updater_plain_factory::create( network_schema_smart_ptr schema, - const std::map& layer_to_dropout_rate_map) const + const std::map& layer_to_dropout_rate_map, + const std::map& layer_to_weight_vector_bound_map) const { - return network_updater_smart_ptr(new network_updater_plain(schema, layer_to_dropout_rate_map, plain_config)); + return network_updater_smart_ptr(new network_updater_plain( + schema, + layer_to_dropout_rate_map, + layer_to_weight_vector_bound_map, + plain_config)); } } } diff --git a/nnforge/plain/network_updater_plain_factory.h b/nnforge/plain/network_updater_plain_factory.h index dd74f2a..4c1e76c 100644 --- a/nnforge/plain/network_updater_plain_factory.h +++ b/nnforge/plain/network_updater_plain_factory.h @@ -32,7 +32,8 @@ namespace nnforge virtual network_updater_smart_ptr create( network_schema_smart_ptr schema, - const std::map& layer_to_dropout_rate_map) const; + const std::map& layer_to_dropout_rate_map, + const std::map& layer_to_weight_vector_bound_map) const; protected: plain_running_configuration_const_smart_ptr plain_config; diff --git a/nnforge/plain/plain.cpp b/nnforge/plain/plain.cpp index ec531c7..5b135a2 100644 --- a/nnforge/plain/plain.cpp +++ b/nnforge/plain/plain.cpp @@ -29,8 +29,6 @@ #include "softmax_layer_tester_plain.h" #include "layer_hessian_plain_factory.h" -#include "layer_updater_plain_factory.h" - #include "absolute_layer_hessian_plain.h" #include "hyperbolic_tangent_layer_hessian_plain.h" #include "average_subsampling_layer_hessian_plain.h" @@ -40,6 +38,7 @@ #include "soft_rectified_linear_layer_hessian_plain.h" #include "softmax_layer_hessian_plain.h" +#include "layer_updater_plain_factory.h" #include "absolute_layer_updater_plain.h" #include "hyperbolic_tangent_layer_updater_plain.h" #include "average_subsampling_layer_updater_plain.h" @@ -49,6 +48,9 @@ #include "soft_rectified_linear_layer_updater_plain.h" #include "softmax_layer_updater_plain.h" +#include "weight_vector_bound_plain_factory.h" +#include "convolution_weight_vector_bound_plain.h" + namespace nnforge { namespace plain @@ -84,6 +86,8 @@ namespace nnforge single_layer_updater_plain_factory::get_mutable_instance().register_layer_updater_plain(layer_updater_plain_smart_ptr(new rectified_linear_layer_updater_plain())); single_layer_updater_plain_factory::get_mutable_instance().register_layer_updater_plain(layer_updater_plain_smart_ptr(new soft_rectified_linear_layer_updater_plain())); single_layer_updater_plain_factory::get_mutable_instance().register_layer_updater_plain(layer_updater_plain_smart_ptr(new softmax_layer_updater_plain())); + + single_weight_vector_bound_factory::get_mutable_instance().register_weight_vector_bound_plain(weight_vector_bound_plain_smart_ptr(new convolution_weight_vector_bound_plain())); } } } diff --git a/nnforge/plain/weight_vector_bound_plain.cpp b/nnforge/plain/weight_vector_bound_plain.cpp new file mode 100644 index 0000000..91fa8ea --- /dev/null +++ b/nnforge/plain/weight_vector_bound_plain.cpp @@ -0,0 +1,31 @@ +/* + * Copyright 2011-2013 Maxim Milakov + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#include "weight_vector_bound_plain.h" + +namespace nnforge +{ + namespace plain + { + weight_vector_bound_plain::weight_vector_bound_plain() + { + } + + weight_vector_bound_plain::~weight_vector_bound_plain() + { + } + } +} diff --git a/nnforge/plain/weight_vector_bound_plain.h b/nnforge/plain/weight_vector_bound_plain.h new file mode 100644 index 0000000..dba8fdc --- /dev/null +++ b/nnforge/plain/weight_vector_bound_plain.h @@ -0,0 +1,56 @@ +/* + * Copyright 2011-2013 Maxim Milakov + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#pragma once + +#include "../layer.h" +#include "plain_running_configuration.h" + +#include "../weight_vector_bound.h" + +#include + +namespace nnforge +{ + namespace plain + { + class weight_vector_bound_plain + { + public: + virtual ~weight_vector_bound_plain(); + + virtual const boost::uuids::uuid& get_uuid() const = 0; + + virtual void normalize_weights( + const weight_vector_bound& bound, + layer_data_list& data, + plain_running_configuration_const_smart_ptr plain_config, + const_layer_smart_ptr layer_schema, + unsigned int updater_count) const = 0; + + protected: + weight_vector_bound_plain(); + + private: + weight_vector_bound_plain(const weight_vector_bound_plain&); + weight_vector_bound_plain& operator =(const weight_vector_bound_plain&); + }; + + typedef std::tr1::shared_ptr weight_vector_bound_plain_smart_ptr; + typedef std::tr1::shared_ptr const_weight_vector_bound_plain_smart_ptr; + typedef std::map weight_vector_bound_map; + } +} diff --git a/nnforge/plain/weight_vector_bound_plain_factory.cpp b/nnforge/plain/weight_vector_bound_plain_factory.cpp new file mode 100644 index 0000000..2617323 --- /dev/null +++ b/nnforge/plain/weight_vector_bound_plain_factory.cpp @@ -0,0 +1,48 @@ +/* + * Copyright 2011-2013 Maxim Milakov + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#include "weight_vector_bound_plain_factory.h" + +#include "../neural_network_exception.h" + +#include +#include + +namespace nnforge +{ + namespace plain + { + bool weight_vector_bound_plain_factory::register_weight_vector_bound_plain(weight_vector_bound_plain_smart_ptr sample_weight_vector_bound_plain) + { + return sample_weight_vector_bound_plain_map.insert(sample_map::value_type(sample_weight_vector_bound_plain->get_uuid(), sample_weight_vector_bound_plain)).second; + } + + bool weight_vector_bound_plain_factory::unregister_weight_vector_bound_plain(const boost::uuids::uuid& layer_guid) + { + return sample_weight_vector_bound_plain_map.erase(layer_guid) == 1; + } + + const_weight_vector_bound_plain_smart_ptr weight_vector_bound_plain_factory::get_updater_plain_layer(const boost::uuids::uuid& layer_guid) const + { + sample_map::const_iterator i = sample_weight_vector_bound_plain_map.find(layer_guid); + + if (i == sample_weight_vector_bound_plain_map.end()) + throw neural_network_exception((boost::format("No plain weight vector bound is registered with id %1%") % layer_guid).str()); + + return i->second; + } + } +} diff --git a/nnforge/plain/weight_vector_bound_plain_factory.h b/nnforge/plain/weight_vector_bound_plain_factory.h new file mode 100644 index 0000000..2e519aa --- /dev/null +++ b/nnforge/plain/weight_vector_bound_plain_factory.h @@ -0,0 +1,48 @@ +/* + * Copyright 2011-2013 Maxim Milakov + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#pragma once + +#include "../layer.h" +#include "weight_vector_bound_plain.h" +#include "plain_running_configuration.h" + +#include +#include +#include +#include + +namespace nnforge +{ + namespace plain + { + class weight_vector_bound_plain_factory + { + public: + bool register_weight_vector_bound_plain(weight_vector_bound_plain_smart_ptr sample_weight_vector_bound_plain); + + bool unregister_weight_vector_bound_plain(const boost::uuids::uuid& layer_guid); + + const_weight_vector_bound_plain_smart_ptr get_updater_plain_layer(const boost::uuids::uuid& layer_guid) const; + + private: + typedef std::map sample_map; + sample_map sample_weight_vector_bound_plain_map; + }; + + typedef boost::serialization::singleton single_weight_vector_bound_factory; + } +} diff --git a/nnforge/weight_vector_bound.cpp b/nnforge/weight_vector_bound.cpp new file mode 100644 index 0000000..a906c81 --- /dev/null +++ b/nnforge/weight_vector_bound.cpp @@ -0,0 +1,29 @@ +/* + * Copyright 2011-2013 Maxim Milakov + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#include "weight_vector_bound.h" + +namespace nnforge +{ + weight_vector_bound::weight_vector_bound(float max_l2_norm) + : max_l2_norm(max_l2_norm) + { + } + + weight_vector_bound::~weight_vector_bound() + { + } +} diff --git a/nnforge/weight_vector_bound.h b/nnforge/weight_vector_bound.h new file mode 100644 index 0000000..52480e8 --- /dev/null +++ b/nnforge/weight_vector_bound.h @@ -0,0 +1,30 @@ +/* + * Copyright 2011-2013 Maxim Milakov + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#pragma once + +namespace nnforge +{ + class weight_vector_bound + { + public: + weight_vector_bound(float max_l2_norm); + + ~weight_vector_bound(); + + float max_l2_norm; + }; +}