bool small_test() {
const int alphabet_size = 5;
const int T = 2;
std::vector<float> activations = {0.1, 0.6, 0.1, 0.1, 0.1,
0.1, 0.1, 0.6, 0.1, 0.1};
// Calculate the score analytically
float expected_score;
{
std::vector<float> probs(activations.size());
softmax(activations.data(), alphabet_size, T, probs.data());
// Score calculation is specific to the given activations above
expected_score = probs[1] * probs[7];
}
std::vector<int> labels = {1, 2};
std::vector<int> label_lengths = {2};
std::vector<int> lengths;
lengths.push_back(T);
float score;
ctcComputeInfo info;
info.loc = CTC_CPU;
info.num_threads = 1;
size_t cpu_alloc_bytes;
throw_on_error(get_workspace_size(label_lengths.data(), lengths.data(),
alphabet_size, lengths.size(), info,
&cpu_alloc_bytes),
"Error: get_workspace_size in small_test");
void* ctc_cpu_workspace = malloc(cpu_alloc_bytes);
throw_on_error(compute_ctc_loss(activations.data(), NULL,
labels.data(), label_lengths.data(),
lengths.data(),
alphabet_size,
lengths.size(),
&score,
ctc_cpu_workspace,
info),
"Error: compute_ctc_loss in small_test");
free(ctc_cpu_workspace);
score = std::exp(-score);
const float eps = 1e-6;
const float lb = expected_score - eps;
const float ub = expected_score + eps;
return (score > lb && score < ub);
}
bool inf_test() {
const int alphabet_size = 15;
const int T = 50;
const int L = 10;
const int minibatch = 1;
std::vector<int> labels = genLabels(alphabet_size, L);
labels[0] = 2;
std::vector<int> label_lengths = {L};
std::vector<float> acts = genActs(alphabet_size * T * minibatch);
for (int i = 0; i < T; ++i)
acts[alphabet_size * i + 2] = -1e30;
std::vector<int> sizes;
sizes.push_back(T);
std::vector<float> grads(alphabet_size * T);
float cost;
ctcComputeInfo info;
info.loc = CTC_CPU;
info.num_threads = 1;
size_t cpu_alloc_bytes;
throw_on_error(get_workspace_size(label_lengths.data(), sizes.data(),
alphabet_size, sizes.size(), info,
&cpu_alloc_bytes),
"Error: get_workspace_size in inf_test");
void* ctc_cpu_workspace = malloc(cpu_alloc_bytes);
throw_on_error(compute_ctc_loss(acts.data(), grads.data(),
labels.data(), label_lengths.data(),
sizes.data(),
alphabet_size,
sizes.size(),
&cost,
ctc_cpu_workspace,
info),
"Error: compute_ctc_loss in inf_test");
free(ctc_cpu_workspace);
bool status = true;
status &= std::isinf(cost);
for (int i = 0; i < alphabet_size * T; ++i)
status &= !std::isnan(grads[i]);
return status;
}
void resize_convolutional_layer(convolutional_layer *l, int w, int h)
{
l->w = w;
l->h = h;
int out_w = convolutional_out_width(*l);
int out_h = convolutional_out_height(*l);
l->out_w = out_w;
l->out_h = out_h;
l->outputs = l->out_h * l->out_w * l->out_c;
l->inputs = l->w * l->h * l->c;
l->output = realloc(l->output, l->batch*l->outputs*sizeof(float));
l->delta = realloc(l->delta, l->batch*l->outputs*sizeof(float));
if(l->batch_normalize){
l->x = realloc(l->x, l->batch*l->outputs*sizeof(float));
l->x_norm = realloc(l->x_norm, l->batch*l->outputs*sizeof(float));
}
#ifdef GPU
cuda_free(l->delta_gpu);
cuda_free(l->output_gpu);
l->delta_gpu = cuda_make_array(l->delta, l->batch*l->outputs);
l->output_gpu = cuda_make_array(l->output, l->batch*l->outputs);
if(l->batch_normalize){
cuda_free(l->x_gpu);
cuda_free(l->x_norm_gpu);
l->x_gpu = cuda_make_array(l->output, l->batch*l->outputs);
l->x_norm_gpu = cuda_make_array(l->output, l->batch*l->outputs);
}
#ifdef CUDNN
cudnn_convolutional_setup(l);
#endif
#endif
l->workspace_size = get_workspace_size(*l);
}
void resize_deconvolutional_layer(layer *l, int h, int w)
{
l->h = h;
l->w = w;
l->out_h = (l->h) * l->stride + l->size/2 - l->pad;
l->out_w = (l->w) * l->stride + l->size/2 - l->pad;
l->outputs = l->out_h * l->out_w * l->out_c;
l->inputs = l->w * l->h * l->c;
l->output = realloc(l->output, l->batch*l->outputs*sizeof(float));
l->delta = realloc(l->delta, l->batch*l->outputs*sizeof(float));
if(l->batch_normalize){
l->x = realloc(l->x, l->batch*l->outputs*sizeof(float));
l->x_norm = realloc(l->x_norm, l->batch*l->outputs*sizeof(float));
}
#ifdef GPU
cuda_free(l->delta_gpu);
cuda_free(l->output_gpu);
l->delta_gpu = cuda_make_array(l->delta, l->batch*l->outputs);
l->output_gpu = cuda_make_array(l->output, l->batch*l->outputs);
if(l->batch_normalize){
cuda_free(l->x_gpu);
cuda_free(l->x_norm_gpu);
l->x_gpu = cuda_make_array(l->output, l->batch*l->outputs);
l->x_norm_gpu = cuda_make_array(l->output, l->batch*l->outputs);
}
#ifdef CUDNN
cudnnSetTensor4dDescriptor(l->dstTensorDesc, CUDNN_TENSOR_NCHW, CUDNN_DATA_FLOAT, l->batch, l->out_c, l->out_h, l->out_w);
cudnnSetTensor4dDescriptor(l->normTensorDesc, CUDNN_TENSOR_NCHW, CUDNN_DATA_FLOAT, 1, l->out_c, 1, 1);
#endif
#endif
l->workspace_size = get_workspace_size(*l);
}
layer make_deconvolutional_layer(int batch, int h, int w, int c, int n, int size, int stride, ACTIVATION activation, int batch_normalize)
{
int i;
layer l = {0};
l.type = DECONVOLUTIONAL;
l.h = h;
l.w = w;
l.c = c;
l.n = n;
l.batch = batch;
l.stride = stride;
l.size = size;
l.weights = calloc(c*n*size*size, sizeof(float));
l.weight_updates = calloc(c*n*size*size, sizeof(float));
l.biases = calloc(n, sizeof(float));
l.bias_updates = calloc(n, sizeof(float));
float scale = 1./sqrt(size*size*c);
for(i = 0; i < c*n*size*size; ++i) l.weights[i] = scale*rand_normal();
for(i = 0; i < n; ++i){
l.biases[i] = scale;
}
l.pad = l.size/2;
l.out_h = (l.h) * l.stride + l.size/2 - l.pad;
l.out_w = (l.w) * l.stride + l.size/2 - l.pad;
l.out_c = n;
l.outputs = l.out_w * l.out_h * l.out_c;
l.inputs = l.w * l.h * l.c;
l.output = calloc(l.batch*l.out_h * l.out_w * n, sizeof(float));
l.delta = calloc(l.batch*l.out_h * l.out_w * n, sizeof(float));
l.forward = forward_deconvolutional_layer;
l.backward = backward_deconvolutional_layer;
l.update = update_deconvolutional_layer;
l.batch_normalize = batch_normalize;
if(batch_normalize){
l.scales = calloc(n, sizeof(float));
l.scale_updates = calloc(n, sizeof(float));
for(i = 0; i < n; ++i){
l.scales[i] = 1;
}
l.mean = calloc(n, sizeof(float));
l.variance = calloc(n, sizeof(float));
l.mean_delta = calloc(n, sizeof(float));
l.variance_delta = calloc(n, sizeof(float));
l.rolling_mean = calloc(n, sizeof(float));
l.rolling_variance = calloc(n, sizeof(float));
l.x = calloc(l.batch*l.outputs, sizeof(float));
l.x_norm = calloc(l.batch*l.outputs, sizeof(float));
}
#ifdef GPU
l.forward_gpu = forward_deconvolutional_layer_gpu;
l.backward_gpu = backward_deconvolutional_layer_gpu;
l.update_gpu = update_deconvolutional_layer_gpu;
if(gpu_index >= 0){
l.weights_gpu = cuda_make_array(l.weights, c*n*size*size);
l.weight_updates_gpu = cuda_make_array(l.weight_updates, c*n*size*size);
l.biases_gpu = cuda_make_array(l.biases, n);
l.bias_updates_gpu = cuda_make_array(l.bias_updates, n);
l.delta_gpu = cuda_make_array(l.delta, l.batch*l.out_h*l.out_w*n);
l.output_gpu = cuda_make_array(l.output, l.batch*l.out_h*l.out_w*n);
if(batch_normalize){
l.mean_gpu = cuda_make_array(l.mean, n);
l.variance_gpu = cuda_make_array(l.variance, n);
l.rolling_mean_gpu = cuda_make_array(l.mean, n);
l.rolling_variance_gpu = cuda_make_array(l.variance, n);
l.mean_delta_gpu = cuda_make_array(l.mean, n);
l.variance_delta_gpu = cuda_make_array(l.variance, n);
l.scales_gpu = cuda_make_array(l.scales, n);
l.scale_updates_gpu = cuda_make_array(l.scale_updates, n);
l.x_gpu = cuda_make_array(l.output, l.batch*l.out_h*l.out_w*n);
l.x_norm_gpu = cuda_make_array(l.output, l.batch*l.out_h*l.out_w*n);
}
}
#ifdef CUDNN
cudnnCreateTensorDescriptor(&l.dstTensorDesc);
cudnnCreateTensorDescriptor(&l.normTensorDesc);
cudnnSetTensor4dDescriptor(l.dstTensorDesc, CUDNN_TENSOR_NCHW, CUDNN_DATA_FLOAT, l.batch, l.out_c, l.out_h, l.out_w);
cudnnSetTensor4dDescriptor(l.normTensorDesc, CUDNN_TENSOR_NCHW, CUDNN_DATA_FLOAT, 1, l.out_c, 1, 1);
#endif
#endif
l.activation = activation;
l.workspace_size = get_workspace_size(l);
fprintf(stderr, "deconv%5d %2d x%2d /%2d %4d x%4d x%4d -> %4d x%4d x%4d\n", n, size, size, stride, w, h, c, l.out_w, l.out_h, l.out_c);
return l;
}
void resize_convolutional_layer(convolutional_layer *l, int w, int h)
{
l->w = w;
l->h = h;
int out_w = convolutional_out_width(*l);
int out_h = convolutional_out_height(*l);
l->out_w = out_w;
l->out_h = out_h;
l->outputs = l->out_h * l->out_w * l->out_c;
l->inputs = l->w * l->h * l->c;
l->output = realloc(l->output,
l->batch*out_h * out_w * l->n*sizeof(float));
l->delta = realloc(l->delta,
l->batch*out_h * out_w * l->n*sizeof(float));
#ifdef GPU
cuda_free(l->delta_gpu);
cuda_free(l->output_gpu);
l->delta_gpu = cuda_make_array(l->delta, l->batch*out_h*out_w*l->n);
l->output_gpu = cuda_make_array(l->output, l->batch*out_h*out_w*l->n);
#ifdef CUDNN
cudnnSetTensor4dDescriptor(l->dsrcTensorDesc, CUDNN_TENSOR_NCHW, CUDNN_DATA_FLOAT, l->batch, l->c, l->h, l->w);
cudnnSetTensor4dDescriptor(l->ddstTensorDesc, CUDNN_TENSOR_NCHW, CUDNN_DATA_FLOAT, l->batch, l->out_c, l->out_h, l->out_w);
cudnnSetFilter4dDescriptor(l->dfilterDesc, CUDNN_DATA_FLOAT, CUDNN_TENSOR_NCHW, l->n, l->c, l->size, l->size);
cudnnSetTensor4dDescriptor(l->srcTensorDesc, CUDNN_TENSOR_NCHW, CUDNN_DATA_FLOAT, l->batch, l->c, l->h, l->w);
cudnnSetTensor4dDescriptor(l->dstTensorDesc, CUDNN_TENSOR_NCHW, CUDNN_DATA_FLOAT, l->batch, l->out_c, l->out_h, l->out_w);
cudnnSetFilter4dDescriptor(l->filterDesc, CUDNN_DATA_FLOAT, CUDNN_TENSOR_NCHW, l->n, l->c, l->size, l->size);
int padding = l->pad ? l->size/2 : 0;
cudnnSetConvolution2dDescriptor(l->convDesc, padding, padding, l->stride, l->stride, 1, 1, CUDNN_CROSS_CORRELATION);
cudnnGetConvolutionForwardAlgorithm(cudnn_handle(),
l->srcTensorDesc,
l->filterDesc,
l->convDesc,
l->dstTensorDesc,
CUDNN_CONVOLUTION_FWD_PREFER_FASTEST,
0,
&l->fw_algo);
cudnnGetConvolutionBackwardDataAlgorithm(cudnn_handle(),
l->filterDesc,
l->ddstTensorDesc,
l->convDesc,
l->dsrcTensorDesc,
CUDNN_CONVOLUTION_BWD_DATA_PREFER_FASTEST,
0,
&l->bd_algo);
cudnnGetConvolutionBackwardFilterAlgorithm(cudnn_handle(),
l->srcTensorDesc,
l->ddstTensorDesc,
l->convDesc,
l->dfilterDesc,
CUDNN_CONVOLUTION_BWD_FILTER_PREFER_FASTEST,
0,
&l->bf_algo);
#endif
#endif
l->workspace_size = get_workspace_size(*l);
}
convolutional_layer make_convolutional_layer(int batch, int h, int w, int c, int n, int size, int stride, int pad, ACTIVATION activation, int batch_normalize, int binary, int xnor)
{
int i;
convolutional_layer l = {0};
l.type = CONVOLUTIONAL;
l.h = h;
l.w = w;
l.c = c;
l.n = n;
l.binary = binary;
l.xnor = xnor;
l.batch = batch;
l.stride = stride;
l.size = size;
l.pad = pad;
l.batch_normalize = batch_normalize;
l.filters = calloc(c*n*size*size, sizeof(float));
l.filter_updates = calloc(c*n*size*size, sizeof(float));
l.biases = calloc(n, sizeof(float));
l.bias_updates = calloc(n, sizeof(float));
// float scale = 1./sqrt(size*size*c);
float scale = sqrt(2./(size*size*c));
for(i = 0; i < c*n*size*size; ++i) l.filters[i] = scale*rand_uniform(-1, 1);
int out_h = convolutional_out_height(l);
int out_w = convolutional_out_width(l);
l.out_h = out_h;
l.out_w = out_w;
l.out_c = n;
l.outputs = l.out_h * l.out_w * l.out_c;
l.inputs = l.w * l.h * l.c;
l.output = calloc(l.batch*out_h * out_w * n, sizeof(float));
l.delta = calloc(l.batch*out_h * out_w * n, sizeof(float));
if(binary){
l.binary_filters = calloc(c*n*size*size, sizeof(float));
l.cfilters = calloc(c*n*size*size, sizeof(char));
l.scales = calloc(n, sizeof(float));
}
if(xnor){
l.binary_filters = calloc(c*n*size*size, sizeof(float));
l.binary_input = calloc(l.inputs*l.batch, sizeof(float));
}
if(batch_normalize){
l.scales = calloc(n, sizeof(float));
l.scale_updates = calloc(n, sizeof(float));
for(i = 0; i < n; ++i){
l.scales[i] = 1;
}
l.mean = calloc(n, sizeof(float));
l.variance = calloc(n, sizeof(float));
l.rolling_mean = calloc(n, sizeof(float));
l.rolling_variance = calloc(n, sizeof(float));
}
#ifdef GPU
l.filters_gpu = cuda_make_array(l.filters, c*n*size*size);
l.filter_updates_gpu = cuda_make_array(l.filter_updates, c*n*size*size);
l.biases_gpu = cuda_make_array(l.biases, n);
l.bias_updates_gpu = cuda_make_array(l.bias_updates, n);
l.scales_gpu = cuda_make_array(l.scales, n);
l.scale_updates_gpu = cuda_make_array(l.scale_updates, n);
l.delta_gpu = cuda_make_array(l.delta, l.batch*out_h*out_w*n);
l.output_gpu = cuda_make_array(l.output, l.batch*out_h*out_w*n);
if(binary){
l.binary_filters_gpu = cuda_make_array(l.filters, c*n*size*size);
}
if(xnor){
l.binary_filters_gpu = cuda_make_array(l.filters, c*n*size*size);
l.binary_input_gpu = cuda_make_array(0, l.inputs*l.batch);
}
if(batch_normalize){
l.mean_gpu = cuda_make_array(l.mean, n);
l.variance_gpu = cuda_make_array(l.variance, n);
l.rolling_mean_gpu = cuda_make_array(l.mean, n);
l.rolling_variance_gpu = cuda_make_array(l.variance, n);
l.mean_delta_gpu = cuda_make_array(l.mean, n);
l.variance_delta_gpu = cuda_make_array(l.variance, n);
l.x_gpu = cuda_make_array(l.output, l.batch*out_h*out_w*n);
l.x_norm_gpu = cuda_make_array(l.output, l.batch*out_h*out_w*n);
}
#ifdef CUDNN
cudnnCreateTensorDescriptor(&l.srcTensorDesc);
cudnnCreateTensorDescriptor(&l.dstTensorDesc);
cudnnCreateFilterDescriptor(&l.filterDesc);
cudnnCreateTensorDescriptor(&l.dsrcTensorDesc);
cudnnCreateTensorDescriptor(&l.ddstTensorDesc);
cudnnCreateFilterDescriptor(&l.dfilterDesc);
cudnnCreateConvolutionDescriptor(&l.convDesc);
cudnnSetTensor4dDescriptor(l.dsrcTensorDesc, CUDNN_TENSOR_NCHW, CUDNN_DATA_FLOAT, l.batch, l.c, l.h, l.w);
cudnnSetTensor4dDescriptor(l.ddstTensorDesc, CUDNN_TENSOR_NCHW, CUDNN_DATA_FLOAT, l.batch, l.out_c, l.out_h, l.out_w);
cudnnSetFilter4dDescriptor(l.dfilterDesc, CUDNN_DATA_FLOAT, CUDNN_TENSOR_NCHW, l.n, l.c, l.size, l.size);
cudnnSetTensor4dDescriptor(l.srcTensorDesc, CUDNN_TENSOR_NCHW, CUDNN_DATA_FLOAT, l.batch, l.c, l.h, l.w);
cudnnSetTensor4dDescriptor(l.dstTensorDesc, CUDNN_TENSOR_NCHW, CUDNN_DATA_FLOAT, l.batch, l.out_c, l.out_h, l.out_w);
cudnnSetFilter4dDescriptor(l.filterDesc, CUDNN_DATA_FLOAT, CUDNN_TENSOR_NCHW, l.n, l.c, l.size, l.size);
int padding = l.pad ? l.size/2 : 0;
cudnnSetConvolution2dDescriptor(l.convDesc, padding, padding, l.stride, l.stride, 1, 1, CUDNN_CROSS_CORRELATION);
cudnnGetConvolutionForwardAlgorithm(cudnn_handle(),
l.srcTensorDesc,
l.filterDesc,
l.convDesc,
l.dstTensorDesc,
CUDNN_CONVOLUTION_FWD_PREFER_FASTEST,
0,
&l.fw_algo);
cudnnGetConvolutionBackwardDataAlgorithm(cudnn_handle(),
l.filterDesc,
l.ddstTensorDesc,
l.convDesc,
l.dsrcTensorDesc,
CUDNN_CONVOLUTION_BWD_DATA_PREFER_FASTEST,
0,
&l.bd_algo);
cudnnGetConvolutionBackwardFilterAlgorithm(cudnn_handle(),
l.srcTensorDesc,
l.ddstTensorDesc,
l.convDesc,
l.dfilterDesc,
CUDNN_CONVOLUTION_BWD_FILTER_PREFER_FASTEST,
0,
&l.bf_algo);
#endif
#endif
l.workspace_size = get_workspace_size(l);
l.activation = activation;
fprintf(stderr, "Convolutional Layer: %d x %d x %d image, %d filters -> %d x %d x %d image\n", h,w,c,n, out_h, out_w, n);
return l;
}
float grad_check(int T, int alphabet_size,
std::vector<float>& acts,
const std::vector<std::vector<int>>& labels,
const std::vector<int>& sizes) {
float epsilon = 1e-2;
const int minibatch = labels.size();
std::vector<int> flat_labels;
std::vector<int> label_lengths;
for (const auto& l : labels) {
flat_labels.insert(flat_labels.end(), l.begin(), l.end());
label_lengths.push_back(l.size());
}
std::vector<float> costs(minibatch);
std::vector<float> grads(acts.size());
ctcComputeInfo info;
info.loc = CTC_CPU;
info.num_threads = 1;
size_t cpu_alloc_bytes;
throw_on_error(get_workspace_size(label_lengths.data(), sizes.data(),
alphabet_size, sizes.size(), info,
&cpu_alloc_bytes),
"Error: get_workspace_size in grad_check");
void* ctc_cpu_workspace = malloc(cpu_alloc_bytes);
throw_on_error(compute_ctc_loss(acts.data(), grads.data(),
flat_labels.data(), label_lengths.data(),
sizes.data(),
alphabet_size,
minibatch,
costs.data(),
ctc_cpu_workspace,
info),
"Error: compute_ctc_loss (0) in grad_check");
float cost = std::accumulate(costs.begin(), costs.end(), 0.);
std::vector<float> num_grad(grads.size());
//perform 2nd order central differencing
for (int i = 0; i < T * alphabet_size * minibatch; ++i) {
std::vector<float> costsP1(minibatch);
std::vector<float> costsP2(minibatch);
acts[i] += epsilon;
throw_on_error(compute_ctc_loss(acts.data(), NULL,
flat_labels.data(), label_lengths.data(),
sizes.data(),
alphabet_size,
minibatch,
costsP1.data(),
ctc_cpu_workspace,
info),
"Error: compute_ctc_loss (1) in grad_check");
acts[i] -= 2 * epsilon;
throw_on_error(compute_ctc_loss(acts.data(), NULL,
flat_labels.data(), label_lengths.data(),
sizes.data(),
alphabet_size,
minibatch,
costsP2.data(),
ctc_cpu_workspace,
info),
"Error: compute_ctc_loss (2) in grad_check");
float costP1 = std::accumulate(costsP1.begin(), costsP1.end(), 0.);
float costP2 = std::accumulate(costsP2.begin(), costsP2.end(), 0.);
acts[i] += epsilon;
num_grad[i] = (costP1 - costP2) / (2 * epsilon);
}
free(ctc_cpu_workspace);
float diff = rel_diff(grads, num_grad);
return diff;
}
convolutional_layer make_convolutional_layer(int batch, int h, int w, int c, int n, int groups, int size, int stride, int padding, ACTIVATION activation, int batch_normalize, int binary, int xnor, int adam)
{
int i;
convolutional_layer l = {0};
l.type = CONVOLUTIONAL;
l.groups = groups;
l.h = h;
l.w = w;
l.c = c;
l.n = n;
l.binary = binary;
l.xnor = xnor;
l.batch = batch;
l.stride = stride;
l.size = size;
l.pad = padding;
l.batch_normalize = batch_normalize;
l.weights = calloc(c/groups*n*size*size, sizeof(float));
l.weight_updates = calloc(c/groups*n*size*size, sizeof(float));
l.biases = calloc(n, sizeof(float));
l.bias_updates = calloc(n, sizeof(float));
l.nweights = c/groups*n*size*size;
l.nbiases = n;
// float scale = 1./sqrt(size*size*c);
float scale = sqrt(2./(size*size*c/l.groups));
//scale = .02;
//for(i = 0; i < c*n*size*size; ++i) l.weights[i] = scale*rand_uniform(-1, 1);
for(i = 0; i < l.nweights; ++i) l.weights[i] = scale*rand_normal();
int out_w = convolutional_out_width(l);
int out_h = convolutional_out_height(l);
l.out_h = out_h;
l.out_w = out_w;
l.out_c = n;
l.outputs = l.out_h * l.out_w * l.out_c;
l.inputs = l.w * l.h * l.c;
l.output = calloc(l.batch*l.outputs, sizeof(float));
l.delta = calloc(l.batch*l.outputs, sizeof(float));
l.forward = forward_convolutional_layer;
l.backward = backward_convolutional_layer;
l.update = update_convolutional_layer;
if(binary){
l.binary_weights = calloc(l.nweights, sizeof(float));
l.cweights = calloc(l.nweights, sizeof(char));
l.scales = calloc(n, sizeof(float));
}
if(xnor){
l.binary_weights = calloc(l.nweights, sizeof(float));
l.binary_input = calloc(l.inputs*l.batch, sizeof(float));
}
if(batch_normalize){
l.scales = calloc(n, sizeof(float));
l.scale_updates = calloc(n, sizeof(float));
for(i = 0; i < n; ++i){
l.scales[i] = 1;
}
l.mean = calloc(n, sizeof(float));
l.variance = calloc(n, sizeof(float));
l.mean_delta = calloc(n, sizeof(float));
l.variance_delta = calloc(n, sizeof(float));
l.rolling_mean = calloc(n, sizeof(float));
l.rolling_variance = calloc(n, sizeof(float));
l.x = calloc(l.batch*l.outputs, sizeof(float));
l.x_norm = calloc(l.batch*l.outputs, sizeof(float));
}
if(adam){
l.m = calloc(l.nweights, sizeof(float));
l.v = calloc(l.nweights, sizeof(float));
l.bias_m = calloc(n, sizeof(float));
l.scale_m = calloc(n, sizeof(float));
l.bias_v = calloc(n, sizeof(float));
l.scale_v = calloc(n, sizeof(float));
}
#ifdef GPU
l.forward_gpu = forward_convolutional_layer_gpu;
l.backward_gpu = backward_convolutional_layer_gpu;
l.update_gpu = update_convolutional_layer_gpu;
if(gpu_index >= 0){
if (adam) {
l.m_gpu = cuda_make_array(l.m, l.nweights);
l.v_gpu = cuda_make_array(l.v, l.nweights);
l.bias_m_gpu = cuda_make_array(l.bias_m, n);
l.bias_v_gpu = cuda_make_array(l.bias_v, n);
l.scale_m_gpu = cuda_make_array(l.scale_m, n);
l.scale_v_gpu = cuda_make_array(l.scale_v, n);
}
l.weights_gpu = cuda_make_array(l.weights, l.nweights);
l.weight_updates_gpu = cuda_make_array(l.weight_updates, l.nweights);
l.biases_gpu = cuda_make_array(l.biases, n);
l.bias_updates_gpu = cuda_make_array(l.bias_updates, n);
l.delta_gpu = cuda_make_array(l.delta, l.batch*out_h*out_w*n);
l.output_gpu = cuda_make_array(l.output, l.batch*out_h*out_w*n);
if(binary){
l.binary_weights_gpu = cuda_make_array(l.weights, l.nweights);
}
if(xnor){
l.binary_weights_gpu = cuda_make_array(l.weights, l.nweights);
l.binary_input_gpu = cuda_make_array(0, l.inputs*l.batch);
}
if(batch_normalize){
l.mean_gpu = cuda_make_array(l.mean, n);
l.variance_gpu = cuda_make_array(l.variance, n);
l.rolling_mean_gpu = cuda_make_array(l.mean, n);
l.rolling_variance_gpu = cuda_make_array(l.variance, n);
l.mean_delta_gpu = cuda_make_array(l.mean, n);
l.variance_delta_gpu = cuda_make_array(l.variance, n);
l.scales_gpu = cuda_make_array(l.scales, n);
l.scale_updates_gpu = cuda_make_array(l.scale_updates, n);
l.x_gpu = cuda_make_array(l.output, l.batch*out_h*out_w*n);
l.x_norm_gpu = cuda_make_array(l.output, l.batch*out_h*out_w*n);
}
#ifdef CUDNN
cudnnCreateTensorDescriptor(&l.normTensorDesc);
cudnnCreateTensorDescriptor(&l.srcTensorDesc);
cudnnCreateTensorDescriptor(&l.dstTensorDesc);
cudnnCreateFilterDescriptor(&l.weightDesc);
cudnnCreateTensorDescriptor(&l.dsrcTensorDesc);
cudnnCreateTensorDescriptor(&l.ddstTensorDesc);
cudnnCreateFilterDescriptor(&l.dweightDesc);
cudnnCreateConvolutionDescriptor(&l.convDesc);
cudnn_convolutional_setup(&l);
#endif
}
#endif
l.workspace_size = get_workspace_size(l);
l.activation = activation;
//fprintf(stderr, "conv %5d %2d x%2d /%2d %4d x%4d x%4d -> %4d x%4d x%4d\n", n, size, size, stride, w, h, c, l.out_w, l.out_h, l.out_c);
return l;
}
int APPLY_SPECIFIC(ctc_cost_cpu)(PyArrayObject * in_activations,
PyArrayObject * in_labels,
PyArrayObject * in_input_lengths,
PyArrayObject ** out_costs,
PyArrayObject ** out_gradients)
{
ctc_context_t ctc_object;
ctc_context_t * context = &ctc_object;
ctc_context_init( context );
if ( !PyArray_IS_C_CONTIGUOUS( in_activations ) )
{
PyErr_SetString( PyExc_RuntimeError,
"ConnectionistTemporalClassification: activations array must be C-contiguous." );
return 1;
}
npy_float32 * activations = (npy_float32 *) PyArray_DATA( in_activations );
create_contiguous_input_lengths( in_input_lengths, &(context->input_lengths) );
if ( NULL == context->input_lengths )
{
// Destroy previous CTC context before returning exception
ctc_context_destroy( context );
PyErr_Format( PyExc_MemoryError,
"ConnectionistTemporalClassification: Could not allocate memory for input lengths" );
return 1;
}
// flatten labels to conform with library memory layout
create_flat_labels( in_labels, &(context->flat_labels), &(context->label_lengths) );
if ( ( NULL == context->label_lengths ) || ( NULL == context->flat_labels ) )
{
// Destroy previous CTC context before returning exception
ctc_context_destroy( context );
PyErr_Format( PyExc_MemoryError,
"ConnectionistTemporalClassification: Could not allocate memory for labels and their lengths" );
return 1;
}
npy_int minibatch_size = PyArray_DIMS( in_activations )[1];
npy_int alphabet_size = PyArray_DIMS( in_activations )[2];
npy_float32 * costs = NULL;
npy_intp cost_size = minibatch_size;
if ( (*out_costs) == NULL || // Symbolic variable has no memory backing
PyArray_NDIM( *out_costs ) != 1 || // or, matrix has the wrong size
PyArray_DIMS( *out_costs )[0] != cost_size )
{
Py_XDECREF( *out_costs );
// Allocate new matrix
*out_costs = (PyArrayObject *) PyArray_ZEROS( 1, &cost_size, NPY_FLOAT32, 0 );
if ( NULL == (*out_costs) )
{
// Destroy previous CTC context before returning exception
ctc_context_destroy( context );
PyErr_Format( PyExc_MemoryError,
"ConnectionistTemporalClassification: Could not allocate memory for CTC costs" );
return 1;
}
}
costs = (npy_float32 *) PyArray_DATA( *out_costs );
npy_float32 * gradients = NULL;
if ( NULL != out_gradients ) // If gradient computation is not disabled
{
if ( NULL == (*out_gradients) || // Symbolic variable has no real backing
PyArray_NDIM( *out_gradients ) != 3 ||
PyArray_DIMS( *out_gradients )[0] != PyArray_DIMS( in_activations )[0] ||
PyArray_DIMS( *out_gradients )[1] != PyArray_DIMS( in_activations )[1] ||
PyArray_DIMS( *out_gradients )[2] != PyArray_DIMS( in_activations )[2] )
{
// Existing matrix is the wrong size. Make a new one.
// Decrement ref counter to existing array
Py_XDECREF( *out_gradients );
// Allocate new array
*out_gradients = (PyArrayObject *) PyArray_ZEROS(3, PyArray_DIMS( in_activations ),
NPY_FLOAT32, 0);
if ( NULL == (*out_gradients) )
{
// Destroy previous CTC context before returning exception
ctc_context_destroy( context );
PyErr_Format( PyExc_MemoryError,
"ConnectionistTemporalClassification: Could not allocate memory for CTC gradients!" );
return 1;
}
}
gradients = (npy_float32 *) PyArray_DATA( *out_gradients );
}
size_t cpu_workspace_size;
int ctc_error;
ctc_error = ctc_check_result( get_workspace_size( context->label_lengths,
context->input_lengths, alphabet_size, minibatch_size, context->options,
&cpu_workspace_size ),
"Failed to obtain CTC workspace size." );
if ( ctc_error ) // Exception is set by ctc_check_result, return error here
{
// Destroy previous CTC context before returning exception
ctc_context_destroy( context );
return 1;
}
context->workspace = malloc( cpu_workspace_size );
if ( NULL == context->workspace )
{
// Destroy previous CTC context before returning exception
ctc_context_destroy( context );
PyErr_Format( PyExc_MemoryError,
"ConnectionistTemporalClassification: Failed to allocate memory for CTC workspace." );
return 1;
}
ctc_error = ctc_check_result( compute_ctc_loss( activations, gradients,
context->flat_labels, context->label_lengths, context->input_lengths,
alphabet_size, minibatch_size, costs, context->workspace,
context->options ), "Failed to compute CTC loss function." );
if ( ctc_error ) // Exception is set by ctc_check_result, return error here
{
ctc_context_destroy( context );
return 1;
}
ctc_context_destroy( context );
return 0;
}
int APPLY_SPECIFIC(ctc_cost_gpu)(PyGpuArrayObject * in_activations,
PyArrayObject * in_labels,
PyArrayObject * in_input_lengths,
PyGpuArrayObject ** out_costs,
PyGpuArrayObject ** out_gradients,
PyGpuContextObject * gpu_context)
{
ctc_context_t ctc_object;
ctc_context_t * context = &ctc_object;
size_t gpu_workspace_size;
int ctc_error = 0;
const size_t num_activations = PyGpuArray_DIMS( in_activations )[0];
const size_t minibatch_size = PyGpuArray_DIMS( in_activations )[1];
const size_t alphabet_size = PyGpuArray_DIMS( in_activations )[2];
const size_t cost_size = minibatch_size;
const size_t grad_dims[3] = { num_activations, minibatch_size, alphabet_size };
float * costs = NULL,
* activations = NULL,
* gradients = NULL;
cuda_enter( gpu_context->ctx );
ctc_context_init( context, gpu_context );
switch (in_activations->ga.typecode)
{
case GA_FLOAT:
activations = (float *) PyGpuArray_DEV_DATA( in_activations );
break;
default:
ctc_context_destroy( context );
cuda_exit( gpu_context->ctx );
PyErr_SetString( PyExc_TypeError,
"GpuConnectionistTemporalClassification: Unsupported type for activations." );
return 1;
}
create_contiguous_input_lengths( in_input_lengths, &(context->input_lengths) );
if ( NULL == context->input_lengths )
{
// Destroy previous CTC context before returning exception
ctc_context_destroy( context );
cuda_exit( gpu_context->ctx );
PyErr_Format( PyExc_MemoryError,
"GpuConnectionistTemporalClassification: Could not allocate memory for input lengths." );
return 1;
}
// flatten labels to conform with library memory layout
create_flat_labels( in_labels, &(context->flat_labels), &(context->label_lengths) );
if ( ( NULL == context->label_lengths ) || ( NULL == context->flat_labels ) )
{
// Destroy previous CTC context before returning exception
ctc_context_destroy( context );
cuda_exit( gpu_context->ctx );
PyErr_Format( PyExc_MemoryError,
"GpuConnectionistTemporalClassification: Could not allocate memory for labels and their lengths." );
return 1;
}
if ( theano_prep_output( out_costs, 1, &cost_size, in_activations->ga.typecode,
GA_C_ORDER, gpu_context ) != 0 )
{
ctc_context_destroy( context );
cuda_exit( gpu_context->ctx );
return 1;
}
GpuArray_memset( &((*out_costs)->ga), 0 );
costs = (float *) PyGpuArray_DEV_DATA( *out_costs );
if ( NULL != out_gradients ) // if gradient computation is not disabled
{
if ( theano_prep_output( out_gradients, 3, grad_dims, in_activations->ga.typecode,
GA_C_ORDER, gpu_context ) != 0 )
{
ctc_context_destroy( context );
cuda_exit( gpu_context->ctx );
return 1;
}
GpuArray_memset( &((*out_gradients)->ga), 0 );
gradients = (float *) PyGpuArray_DEV_DATA( *out_gradients );
}
ctc_error = ctc_check_result( get_workspace_size( context->label_lengths,
context->input_lengths, alphabet_size, minibatch_size, context->options,
&gpu_workspace_size ),
"Failed to obtain CTC workspace size." );
if ( ctc_error ) // Exception is set by ctc_check_result, return error here
{
// Destroy previous CTC context before returning exception
ctc_context_destroy( context );
cuda_exit( gpu_context->ctx );
return 1;
}
context->workspace = gpudata_alloc( gpu_context->ctx, gpu_workspace_size, NULL, 0, NULL );
if ( NULL == context->workspace )
{
ctc_context_destroy( context );
cuda_exit( gpu_context->ctx );
PyErr_Format( PyExc_MemoryError,
"GpuConnectionistTemporalClassification: Failed to allocate memory for CTC workspace." );
return 1;
}
cuda_wait( in_activations->ga.data, GPUARRAY_CUDA_WAIT_READ );
cuda_wait( (*out_costs)->ga.data, GPUARRAY_CUDA_WAIT_WRITE );
if ( out_gradients != NULL )
cuda_wait( (*out_gradients)->ga.data, GPUARRAY_CUDA_WAIT_WRITE );
ctc_error = ctc_check_result( compute_ctc_loss( activations, gradients,
context->flat_labels, context->label_lengths, context->input_lengths,
alphabet_size, minibatch_size, costs, *(void **)context->workspace,
context->options ), "Failed to compute CTC loss function." );
cuda_record( in_activations->ga.data, GPUARRAY_CUDA_WAIT_READ );
cuda_record( (*out_costs)->ga.data, GPUARRAY_CUDA_WAIT_WRITE );
if ( out_gradients != NULL )
cuda_record( (*out_gradients)->ga.data, GPUARRAY_CUDA_WAIT_WRITE );
if ( ctc_error ) // Exception is set by ctc_check_result, return error here
{
ctc_context_destroy( context );
cuda_exit( gpu_context->ctx );
return 1;
}
ctc_context_destroy( context );
cuda_exit( gpu_context->ctx );
return 0;
}
bool options_test() {
const int alphabet_size = 6;
const int T = 5;
const int minibatch = 2;
std::vector<float> activations =
{0.633766, 0.221185, 0.0917319, 0.0129757, 0.0142857, 0.0260553,
0.30176, 0.28562, 0.0831517, 0.0862751, 0.0816851, 0.161508,
0.111121, 0.588392, 0.278779, 0.0055756, 0.00569609, 0.010436,
0.24082, 0.397533, 0.0557226, 0.0546814, 0.0557528, 0.19549,
0.0357786, 0.633813, 0.321418, 0.00249248, 0.00272882, 0.0037688,
0.230246, 0.450868, 0.0389607, 0.038309, 0.0391602, 0.202456,
0.0663296, 0.643849, 0.280111, 0.00283995, 0.0035545, 0.00331533,
0.280884, 0.429522, 0.0326593, 0.0339046, 0.0326856, 0.190345,
0.458235, 0.396634, 0.123377, 0.00648837, 0.00903441, 0.00623107,
0.423286, 0.315517, 0.0338439, 0.0393744, 0.0339315, 0.154046};
std::vector<float> expected_grads = // from tensorflow
{-0.366234, 0.221185, 0.0917319, 0.0129757, 0.0142857, 0.0260553,
-0.69824, 0.28562, 0.0831517, 0.0862751, 0.0816851, 0.161508,
0.111121, -0.411608, 0.278779, 0.0055756, 0.00569609, 0.010436,
0.24082, -0.602467, 0.0557226, 0.0546814, 0.0557528, 0.19549,
0.0357786, 0.633813, -0.678582, 0.00249248, 0.00272882, 0.0037688,
0.230246, 0.450868, 0.0389607, 0.038309, 0.0391602, -0.797544,
0.0663296, -0.356151, 0.280111, 0.00283995, 0.0035545, 0.00331533,
0.280884, -0.570478, 0.0326593, 0.0339046, 0.0326856, 0.190345,
-0.541765, 0.396634, 0.123377, 0.00648837, 0.00903441, 0.00623107,
-0.576714, 0.315517, 0.0338439, 0.0393744, 0.0339315, 0.154046};
// Calculate the expected scores analytically
std::vector<double> expected_scores(2);
auto& a = activations;
expected_scores[0] =
-std::log(a[offset(0, 0, 0)] * a[offset(1, 0, 1)] * a[offset(2, 0, 2)]
* a[offset(3, 0, 1)] * a[offset(4, 0, 0)]);
expected_scores[1] = 5.42262; // from tensorflow
// now take the log to account for the softmax
for (auto& a : activations) {
a = std::log(a);
}
std::vector<int> labels = {0, 1, 2, 1, 0,
0, 1, 1, 0};
std::vector<int> label_lengths = {5, 4};
std::vector<int> lengths = {5, 5};
std::vector<float> grads(alphabet_size * T * minibatch);
std::vector<float> scores(2);
ctcOptions options{};
options.loc = CTC_CPU;
options.num_threads = 1;
options.blank_label = 5;
size_t cpu_alloc_bytes;
throw_on_error(get_workspace_size(label_lengths.data(), lengths.data(),
alphabet_size, lengths.size(), options,
&cpu_alloc_bytes),
"Error: get_workspace_size in options_test");
void* ctc_cpu_workspace = malloc(cpu_alloc_bytes);
throw_on_error(compute_ctc_loss(activations.data(), grads.data(),
labels.data(), label_lengths.data(),
lengths.data(),
alphabet_size,
lengths.size(),
scores.data(),
ctc_cpu_workspace,
options),
"Error: compute_ctc_loss in options_test");
free(ctc_cpu_workspace);
const double eps = 1e-4;
bool result = true;
for (int i = 0; i < grads.size(); i++) {
const double lb = expected_grads[i] - eps;
const double ub = expected_grads[i] + eps;
if (!(grads[i] > lb && grads[i] < ub)) {
std::cerr << "grad mismatch in options_test"
<< " expected grad: " << expected_grads[i]
<< " calculated score: " << grads[i]
<< " !(" << lb << " < " << grads[i]
<< " < " << ub << ")" << std::endl;
result = false;
}
}
for (int i = 0; i < 2; i++) {
const double lb = expected_scores[i] - eps;
const double ub = expected_scores[i] + eps;
if (!(scores[i] > lb && scores[i] < ub)) {
std::cerr << "score mismatch in options_test"
<< " expected score: " << expected_scores[i]
<< " calculated score: " << scores[i]
<< " !(" << lb << " < " << scores[i]
<< " < " << ub << ")" << std::endl;
result = false;
}
}
return result;
}