/* CUDA_VALUE.H: headers for CUDA kernels and common definitions
 *
 * Augustus Lidaka and Shitanshu Aggarwal
 *
 * Created 2009 July 1
 * Modified 2010 June 30
 */
#ifndef __CUDA_VALUE_H__
#define __CUDA_VALUE_H__

// Maximum number of blocks
#define MAX_NUM_BLOCKS 65535

// Maximum number of threads per block
#define MAX_THREADS_PER_BLOCK 512

// Error codes for parmarginals
#define UINT_MAX_IS_NAN 1
#define FLOAT_MAX_IS_NAN 2
#define EXP_Z_SANS_MAX_IS_NAN 3
#define LOG_Z_SANS_MAX_IS_NAN 4
#define LOG_Z_IS_NAN 5
#define LOG_B_IS_NAN 6
#define EXP_Z_SANS_MAX_IS_ZERO 7
#define EXP_Z_SANS_MAX_IS_INF 8
#define EXP_Z_SANS_MAX_IS_NEG 9
#define THREAD_Z_SANS_MAX_IS_NEG 10
#define THREAD_Z_SANS_MAX_IS_INF 11


/***************************************************************************
 * PARMARGINALS: calculate the marginals (in place) and check for
 *    problematic values
 *
 *
 * Purpose:
 *
 *    Calculate the marginals from the energy. Also store the logprob of the
 *    true label for each instance in a new vector for convenience and
 *    efficiency (rather than indexing into B many times later).
 *
 *
 * Preconditions:
 *
 *    At least as many blocks (gridDim.x) must be spawned as N, the number
 *    of instances, but N must be fewer than the maximum number of blocks
 *    possible on a CUDA device.
 *
 *    N < 65536. Each block takes exactly one instance and the maximum
 *    size of the block grid along any dimension is 65536.
 *
 *
 * Postconditions:
 * 
 *    If any values in B are [-]Inf or NaN, the error flag is immediately
 *    set and we try to do as little further computation as possible, i.e.,
 *    return as quickly as possible to begin a new iteration.
 *
 *
 * Input Parameters:
 *
 *    float *B              Energy matrix which is set to the marginals. L*N
 *                          in Matlab terms; i.e., indices 0:(L-1) represent
 *                          values for the first instance.
 *    const unsigned int *Y Length N vector containing true labels of each
 *                          instance B.
 *    const float *I        Length N vector of instance weights.
 *    unsigned int N        Length (number of columns) of B. Number of
 *                          training instances.
 *    unsigned int L        Height (number of rows) of B. Number of
 *                          training labels.
 *    float eps             Epsilon window around zero (from Matlab); used
 *                          in testing for problematic values in B.
 *    float *vv             N-element vector which will contain the
 *                          likelihood of true labels for each instance.
 *                          Each kernel will modify only one entry (kernel
 *                          with index idx should modify only vv[idx]).
 *    bool *error           Pointer to a bool which is false, but changed to
 *                          true if there are any problematic values in vv.
 *
 *
 * Output (Side Effects):
 *
 *    float *B              Modifies B in place; contains marginals of
 *                          presumed instances and weights from another part
 *                          of value.
 *    float *vv             Likelihood of true labels for each instance.
 *                          Each kernel will modify only one entry (kernel
 *                          with index idx should modify only vv[idx]).
 *    bool *error           If we detect any infinite, NaN, or positive
 *                          values in the likelihood of B, something went
 *                          wrong and this will be set to true. Otherwise,
 *                          false.
 *
 *
 * Philosophy:
 *
 *    Each block has three shared variables: uarray[MAX_THREADS_PER_BLOCK],
 *    Z[MAX_THREADS_PER_BLOCK], and fmx. uarray is an array of uints used in
 *    conjunction with the float/uint bijection (described in parmax.h) and
 *    CUDA atomic functions to find the max of a section of B; fmx holds the
 *    float value of that maximum; and Z is an array of floats used in
 *    summing across intermediate values for normalization. Computation of
 *    marginals then continues in parallel. If we encounter any problematic
 *    values in B, set an error flag and terminate as soon as possible. If
 *    no error is detected, B = exp(B) * I (scale each row of B by the
 *    respective entry in I), set vv, and return.
 *
 **************************************************************************/
__global__ void parmarginals(float *B, const unsigned int *Y, const float
    *I, unsigned int N, int L, float eps, float *vv, bool *error);


#endif



/*

  Copyright 2009,2010 Augustus Lidaka, Shitanshu Aggarwal, Jerod Weinman

  This file is part of CUDAMAXENT.

    CUDAMAXENT is free software: you can redistribute it and/or modify
    it under the terms of the GNU General Public License as published by
    the Free Software Foundation, either version 3 of the License, or
    (at your option) any later version.

    CUDAMAXENT is distributed in the hope that it will be useful, but
    WITHOUT ANY WARRANTY; without even the implied warranty of
    MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the GNU
    General Public License for more details.

    You should have received a copy of the GNU General Public License
    along with CUDAMAXENT.  If not, see <http://www.gnu.org/licenses/>.

*/
