<commit_before>/* backprop_cuda.cc                                                -*- C++ -*-
   <NAME>, 25 May 2009
   Copyright (c) 2009 <NAME>.  All rights reserved.

   CUDA-based backprop implementation.
*/


#include "arch/exception.h"
#include "compiler/compiler.h"
#include <cstdio>
#include <iostream>
#include <boost/timer.hpp>
#include <boost/utility.hpp>
#include "arch/cuda/device_data.h"
#include "arch/cuda/atomic.h"
#include "math/xdiv.h"
#include "perceptron_defs.h"
#include <vector>
#include "backprop_cuda.h"
#include "fixed_point_accum.h"

using namespace std;


/* TODO:
   - Allow it to run with max_width > 512 (maximum thread block width)
   - tanh function that gives bit-for-bit equivalent results as on the
     host
   - Remove learning rate from the update (apply it when updating the weights)
     and use a constant that conditions the numbers to work well within the
     range of the update
   - Try using textures for the W arrays (caching could make a big difference)
*/


typedef ML::FixedPointAccum32 UpdateFloat;
//typedef float UpdateFloat;

/** Given an activation function and an input, apply that activation
    function */
__device__ float transform(float input, int activation)
{
    switch (activation) {
    case ML::ACT_TANH: {
        return tanh(input);
        //float exp2i = __expf(input + input);
        //return __fdividef(exp2i - 1.0f, exp2i + 1.0f);
    }
    case ML::ACT_IDENTITY: return input;
    default:
        return 0.0;
    }
}

/** Given an output and an error, what's the delta (derivative * error)? */
__device__ float delta(float output, float error, int activation)
{
    switch (activation) {
    case ML::ACT_TANH:
        return (1.0f - output * output) * error;
    case ML::ACT_IDENTITY: return output * error; 
    default:
        return 0.0;
    }
}

texture<float, 1, cudaReadModeElementType> weights_tex;
texture<float, 1, cudaReadModeElementType> biases_tex;;


#if !defined(__DEVICE_EMULATION__)

template<const texture<float, 1, cudaReadModeElementType> & Tex>
struct WeightsAccess {
    const float * base;  // if zero, then texture access
    int offset;

    __device__ WeightsAccess(const float * base = 0)
        : base(base), offset(0)
    {
    }

    __device__ void init(const float * base)
    {
        this->base = base;
        offset = 0;
    }

    __device__ void operator += (int val)
    {
        offset += val;
    }

    __device__ void operator -= (int val)
    {
        offset -= val;
    }

    __device__ float operator [] (int ofs)
    {
        if (base) return base[offset + ofs];
        else return tex1Dfetch(Tex, offset + ofs);
    }
};

#endif

/** Train a fully-connected neural network architecture via backpropagation
    one a single training example.  The work is split over all of the cores
    within a single multiprocessor.  (So, on a Geforce 260 core 216, we have
    27 multiprocessors with 8 cores each, and so we could train on 27 different
    feature vectors in parallel.
*/

#define N 4
#define train_N_examples train_4_examples
#include "backprop_cuda_train_N_examples.cu"
#undef N
#undef train_N_examples

#define N 1
#define train_N_examples train_1_example
#include "backprop_cuda_train_N_examples.cu"
#undef N
#undef train_N_examples

#include "backprop_cuda_one_example.cu"

__global__ void
train_examples_kernel(const float * feature_vectors,  // feature vector [ni]
                      int feature_vector_width,
                      const int * labels,
                      const float * example_weights,
                      int num_layers,
                      const float * w,  // weights for each layer
                      const float * biases, // for each layer
                      const int * architecture,
                      const int * w_strides,
                      UpdateFloat * const * w_updates, // wt updates for each layer
                      UpdateFloat * const * b_updates, // bias upd for each layer
                      int activation,            // activation function
                      float inhibit, // target value for inhibited neuron)
                      float fire,   // target value for firing neuron
                      float learning_rate,
                      int num_threads_in_block,
                      int num_threads_on_multiprocessor,
                      int total_neurons,
                      float * layer_outputs,  // scratch space[total neurons]
                      int examples_per_block,
                      int total_num_examples,
                      int max_width,
                      bool use_textures)
{
    const unsigned block_num  = blockIdx.x;
    
    /* Where we accumulate our errors, layer by layer.  The size is that of
       the largest dimension. */
    extern __shared__ float scratch[];
    
    /* The layer outputs (activation of the neurons).  This is where the
       shared memory goes to.  Note that we store only the activated outputs,
       not the inputs.

       blockDim.x gives us the number of threads, which is also the size of
       the errors array, so that our layer outputs have to start at this
       offset.
    */

    // Get our private scratch memory for this block
    layer_outputs += block_num * total_neurons;
    
    unsigned example_num_base = block_num * examples_per_block;
    unsigned last_example = min(total_num_examples, example_num_base + examples_per_block);

    unsigned example_num = example_num_base;

#if defined(__DEVICE_EMULATION__)
    const float * weights_access = w;
    const float * biases_access = biases;
#else
    WeightsAccess<weights_tex> weights_access;
    WeightsAccess<biases_tex> biases_access;

    if (!use_textures) {
        weights_access.init(w);
        biases_access.init(biases);
    }
#endif

#if 0
    for (;  example_num < last_example;  example_num += N) {

        const float * input = feature_vectors + example_num * feature_vector_width;
        train_4_examples(input,
                         labels + example_num,
                         example_weights + example_num,
                         last_example - example_num,
                         num_layers, scratch,
                         weights_access, biases_access,
                         architecture, w_strides,
                         w_updates, b_updates,
                         activation, inhibit, fire, learning_rate,
                         num_threads_in_block,
                         num_threads_on_multiprocessor,
                         total_neurons, max_width, layer_outputs);
    }
#elif 1
    // Do any others singly
    for (;  example_num < last_example;  ++example_num) {

        const float * input
            = feature_vectors + example_num * feature_vector_width;

        train_4_examples(input,
                         labels + example_num,
                         example_weights + example_num,
                         1 /* num valid examples */,<commit_msg>Fixed improper num example calcuation<commit_after>/* backprop_cuda.cc                                                -*- C++ -*-
   <NAME>, 25 May 2009
   Copyright (c) 2009 <NAME>.  All rights reserved.

   CUDA-based backprop implementation.
*/


#include "arch/exception.h"
#include "compiler/compiler.h"
#include <cstdio>
#include <iostream>
#include <boost/timer.hpp>
#include <boost/utility.hpp>
#include "arch/cuda/device_data.h"
#include "arch/cuda/atomic.h"
#include "math/xdiv.h"
#include "perceptron_defs.h"
#include <vector>
#include "backprop_cuda.h"
#include "fixed_point_accum.h"

using namespace std;


/* TODO:
   - Allow it to run with max_width > 512 (maximum thread block width)
   - tanh function that gives bit-for-bit equivalent results as on the
     host
   - Remove learning rate from the update (apply it when updating the weights)
     and use a constant that conditions the numbers to work well within the
     range of the update
   - Try using textures for the W arrays (caching could make a big difference)
*/


typedef ML::FixedPointAccum32 UpdateFloat;
//typedef float UpdateFloat;

/** Given an activation function and an input, apply that activation
    function */
__device__ float transform(float input, int activation)
{
    switch (activation) {
    case ML::ACT_TANH: {
        return tanh(input);
        //float exp2i = __expf(input + input);
        //return __fdividef(exp2i - 1.0f, exp2i + 1.0f);
    }
    case ML::ACT_IDENTITY: return input;
    default:
        return 0.0;
    }
}

/** Given an output and an error, what's the delta (derivative * error)? */
__device__ float delta(float output, float error, int activation)
{
    switch (activation) {
    case ML::ACT_TANH:
        return (1.0f - output * output) * error;
    case ML::ACT_IDENTITY: return output * error; 
    default:
        return 0.0;
    }
}

texture<float, 1, cudaReadModeElementType> weights_tex;
texture<float, 1, cudaReadModeElementType> biases_tex;;


#if !defined(__DEVICE_EMULATION__)

template<const texture<float, 1, cudaReadModeElementType> & Tex>
struct WeightsAccess {
    const float * base;  // if zero, then texture access
    int offset;

    __device__ WeightsAccess(const float * base = 0)
        : base(base), offset(0)
    {
    }

    __device__ void init(const float * base)
    {
        this->base = base;
        offset = 0;
    }

    __device__ void operator += (int val)
    {
        offset += val;
    }

    __device__ void operator -= (int val)
    {
        offset -= val;
    }

    __device__ float operator [] (int ofs)
    {
        if (base) return base[offset + ofs];
        else return tex1Dfetch(Tex, offset + ofs);
    }
};

#endif

/** Train a fully-connected neural network architecture via backpropagation
    one a single training example.  The work is split over all of the cores
    within a single multiprocessor.  (So, on a Geforce 260 core 216, we have
    27 multiprocessors with 8 cores each, and so we could train on 27 different
    feature vectors in parallel.
*/

#define N 4
#define train_N_examples train_4_examples
#include "backprop_cuda_train_N_examples.cu"
#undef N
#undef train_N_examples

#define N 1
#define train_N_examples train_1_example
#include "backprop_cuda_train_N_examples.cu"
#undef N
#undef train_N_examples

#include "backprop_cuda_one_example.cu"

__global__ void
train_examples_kernel(const float * feature_vectors,  // feature vector [ni]
                      int feature_vector_width,
                      const int * labels,
                      const float * example_weights,
                      int num_layers,
                      const float * w,  // weights for each layer
                      const float * biases, // for each layer
                      const int * architecture,
                      const int * w_strides,
                      UpdateFloat * const * w_updates, // wt updates for each layer
                      UpdateFloat * const * b_updates, // bias upd for each layer
                      int activation,            // activation function
                      float inhibit, // target value for inhibited neuron)
                      float fire,   // target value for firing neuron
                      float learning_rate,
                      int num_threads_in_block,
                      int num_threads_on_multiprocessor,
                      int total_neurons,
                      float * layer_outputs,  // scratch space[total neurons]
                      int examples_per_block,
                      int total_num_examples,
                      int max_width,
                      bool use_textures)
{
    const unsigned block_num  = blockIdx.x;
    
    /* Where we accumulate our errors, layer by layer.  The size is that of
       the largest dimension. */
    extern __shared__ float scratch[];
    
    /* The layer outputs (activation of the neurons).  This is where the
       shared memory goes to.  Note that we store only the activated outputs,
       not the inputs.

       blockDim.x gives us the number of threads, which is also the size of
       the errors array, so that our layer outputs have to start at this
       offset.
    */

    // Get our private scratch memory for this block
    layer_outputs += block_num * total_neurons;
    
    unsigned example_num_base = block_num * examples_per_block;
    unsigned last_example = min(total_num_examples, example_num_base + examples_per_block);

    unsigned example_num = example_num_base;

#if defined(__DEVICE_EMULATION__)
    const float * weights_access = w;
    const float * biases_access = biases;
#else
    WeightsAccess<weights_tex> weights_access;
    WeightsAccess<biases_tex> biases_access;

    if (!use_textures) {
        weights_access.init(w);
        biases_access.init(biases);
    }
#endif

#if 0
    for (;  example_num < last_example;  example_num += 4) {

        const float * input = feature_vectors + example_num * feature_vector_width;
        train_4_examples(input,
                         labels + example_num,
                         example_weights + example_num,
                         min(4, last_example - example_num),
                         num_layers, scratch,
                         weights_access, biases_access,
                         architecture, w_strides,
                         w_updates, b_updates,
                         activation, inhibit, fire, learning_rate,
                         num_threads_in_block,
                         num_threads_on_multiprocessor,
                         total_neurons, max_width, layer_outputs);
    }
#elif 1
    // Do any others singly
    for (;  example_num < last_example;  ++example_num) {

        const float * input
            = feature_vectors + example_num * feature_vector_width;

        train_4_examples(input,
                         labels + example_num,
                         example_weights + example_num,
                         1 /* num valid examples */,
{'fraction_non_alphanumeric': 0.07462153670379891, 'fraction_numerical': 0.007997714938588975, 'mean_word_length': 2.0861613045394445, 'pattern_counts': {'":': 0, '<': 33, '<?xml version=': 0, '>': 33, 'https://': 0, 'lorem ipsum': 0, 'www.': 0, 'xml': 0}, 'pii_count': 0, 'substrings_counts': 0, 'word_list_counts': {'cursed_substrings.json': 4, 'profanity_word_list.json': 0, 'sexual_word_list.json': 0, 'zh_pornsignals.json': 0}}
{'dir': 'git-commits-cleaned', 'id': '4985674', 'n_tokens_mistral': 4065, 'n_tokens_neox': 3781, 'n_words': 1241}