FM / core

/usr/local/lib64/python3.6/site-packages/torch/include/ATen/core/

/usr/local/lib64/python3.6/site-packages/torch/include/ATen/core

Name	Size	Mode	Actions
boxing/	-	0755	rm
dispatch/	-	0755	rm
op_registration/	-	0755	rm
alias_info.h	2986	0644	edit dl rm
Array.h	768	0644	edit dl rm
ATenGeneral.h	45	0644	edit dl rm
ATenOpList.h	246	0644	edit dl rm
aten_interned_strings.h	25389	0644	edit dl rm
Backtrace.h	59	0644	edit dl rm
blob.h	5422	0644	edit dl rm
builtin_function.h	3649	0644	edit dl rm
DeprecatedTypeProperties.h	3773	0644	edit dl rm
DeprecatedTypePropertiesRegistry.h	795	0644	edit dl rm
Dict.h	13195	0644	edit dl rm
Dict_inl.h	7996	0644	edit dl rm
Dimname.h	1188	0644	edit dl rm
DimVector.h	247	0644	edit dl rm
DistributionsHelper.h	12594	0644	edit dl rm
Formatting.h	959	0644	edit dl rm
function.h	2145	0644	edit dl rm
functional.h	1460	0644	edit dl rm
function_schema.h	13577	0644	edit dl rm
function_schema_inl.h	9319	0644	edit dl rm
Generator.h	4935	0644	edit dl rm
grad_mode.h	210	0644	edit dl rm
interned_strings.h	25332	0644	edit dl rm
interned_strings_class.h	770	0644	edit dl rm
ivalue.h	38823	0644	edit dl rm
ivalue_inl.h	59963	0644	edit dl rm
ivalue_to.h	756	0644	edit dl rm
jit_type.h	75971	0644	edit dl rm
jit_type_base.h	6508	0644	edit dl rm
LegacyTypeDispatch.h	4626	0644	edit dl rm
List.h	15667	0644	edit dl rm
List_inl.h	11012	0644	edit dl rm
Macros.h	44	0644	edit dl rm
MT19937RNGEngine.h	6410	0644	edit dl rm
NamedTensor.h	5050	0644	edit dl rm
operator_name.h	3018	0644	edit dl rm
PhiloxRNGEngine.h	6496	0644	edit dl rm
PythonModeTLS.h	403	0644	edit dl rm
qualified_name.h	4358	0644	edit dl rm
QuantizerBase.h	2443	0644	edit dl rm
Range.h	418	0644	edit dl rm
Reduction.h	461	0644	edit dl rm
rref_interface.h	1144	0644	edit dl rm
Scalar.h	29	0644	edit dl rm
ScalarType.h	33	0644	edit dl rm
stack.h	6034	0644	edit dl rm
Tensor.h	1756	0644	edit dl rm
TensorAccessor.h	10296	0644	edit dl rm
TensorBase.h	32767	0644	edit dl rm
TensorBody.h	247555	0644	edit dl rm
TransformationHelper.h	6911	0644	edit dl rm
typeid.h	29	0644	edit dl rm
UndefinedTensorImpl.h	42	0644	edit dl rm
UnsafeFromTH.h	708	0644	edit dl rm
VariableHooksInterface.h	3312	0644	edit dl rm
Variadic.h	2257	0644	edit dl rm
Vitals.h	2305	0644	edit dl rm

Edit: /usr/local/lib64/python3.6/site-packages/torch/include/ATen/core/PhiloxRNGEngine.h (6496B)

#pragma once

// define constants like M_PI and C keywords for MSVC
#ifdef _MSC_VER
#define _USE_MATH_DEFINES
#include <math.h>
#endif

#include <stdint.h>

#ifdef __CUDACC__
#include <cuda.h>
#endif

#include <ATen/core/Array.h>
#include <c10/macros/Macros.h>
#include <c10/util/Exception.h>
#include <c10/util/Half.h>
#include <cmath>

namespace at {

// typedefs for holding vector data
namespace detail {

typedef at::detail::Array<uint32_t, 4> UINT4;
typedef at::detail::Array<uint32_t, 2> UINT2;
typedef at::detail::Array<double, 2> DOUBLE2;
typedef at::detail::Array<float, 2> FLOAT2;

} // namespace detail

/**
 * Note [Philox Engine implementation]
 * ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
 * Originally implemented in PyTorch's fusion compiler
 * Refer to: http://www.thesalmons.org/john/random123/papers/random123sc11.pdf
 * for details regarding the engine.
 *
 * Note that currently this implementation of the philox engine is not used
 * anywhere except for tests in cpu_generator_test.cpp. However, this engine
 * will replace curandStatePhilox4_32_10_t in the future.
 *
 * The philox engine takes a seed value, a subsequeunce
 * for starting the generation and an offset for the subsequence.
 * Think of this engine as an algorithm producing a huge array. We are
 * parallelizing this array by partitioning the huge array and assigning
 * a thread index to each partition. In other words, each seed value
 * (there are 2^64 possible seed values) gives a sub array of size
 * 2^128 (each element in that array is a 128 bit number). Reasoning
 * behind the array being of size 2^128 is, there are 2^64 possible
 * thread index value and there is an array of size 2^64 for each of
 * those thread index. Hence 2^64 * 2^64 = 2^128 for each seed value.
 *
 * In short, this generator can produce 2^64 (seed values) * 2^128 (number
 * of elements in an array given by a seed value) = 2^192 values.
 *
 * Arguments:
 * seed:        Seed values could be any number from 0 to 2^64-1.
 * subsequence: Subsequence is just the cuda thread indexing with:
 *              - blockIdx.x * blockDim.x + threadIdx.x
 * offset:      The offset variable in PhiloxEngine  decides how many 128-bit
 *              random numbers to skip (i.e. how many groups of 4, 32-bit numbers to skip)
 *              and hence really decides the total number of randoms that can be achieved
 *              for the given subsequence.
 */

class philox_engine {
public:

C10_HOST_DEVICE inline explicit philox_engine(uint64_t seed = 67280421310721,
                                 uint64_t subsequence = 0,
                                 uint64_t offset = 0) {
    key[0] = static_cast<uint32_t>(seed);
    key[1] = static_cast<uint32_t>(seed >> 32);
    counter = detail::UINT4(0);
    counter[2] = static_cast<uint32_t>(subsequence);
    counter[3] = static_cast<uint32_t>(subsequence >> 32);
    STATE = 0;
    incr_n(offset);
  }

/**
   * Produces a unique 32-bit pseudo random number on every invocation
   */
  C10_HOST_DEVICE inline uint32_t operator()() {
    if(STATE == 0) {
      detail::UINT4 counter_ = counter;
      detail::UINT2 key_ = key;

counter_ = single_round(counter_, key_);
      key_[0] += (kPhilox10A); key_[1] += (kPhilox10B);
      counter_ = single_round(counter_, key_);
      key_[0] += (kPhilox10A); key_[1] += (kPhilox10B);
      counter_ = single_round(counter_, key_);
      key_[0] += (kPhilox10A); key_[1] += (kPhilox10B);
      counter_ = single_round(counter_, key_);
      key_[0] += (kPhilox10A); key_[1] += (kPhilox10B);
      counter_ = single_round(counter_, key_);
      key_[0] += (kPhilox10A); key_[1] += (kPhilox10B);
      counter_ = single_round(counter_, key_);
      key_[0] += (kPhilox10A); key_[1] += (kPhilox10B);
      counter_ = single_round(counter_, key_);
      key_[0] += (kPhilox10A); key_[1] += (kPhilox10B);
      counter_ = single_round(counter_, key_);
      key_[0] += (kPhilox10A); key_[1] += (kPhilox10B);
      counter_ = single_round(counter_, key_);
      key_[0] += (kPhilox10A); key_[1] += (kPhilox10B);

output = single_round(counter_, key_);
      incr();
    }
    uint32_t ret = output[STATE];
    STATE = (STATE + 1) & 3;
    return ret;
  }

/**
   * Function that Skips N 128 bit numbers in a subsequence
   */
  C10_HOST_DEVICE inline void incr_n(uint64_t n) {
    uint32_t nlo = static_cast<uint32_t>(n);
    uint32_t nhi = static_cast<uint32_t>(n >> 32);
    counter[0] += nlo;
    // if overflow in x has occurred, carry over to nhi
    if (counter[0] < nlo) {
      nhi++;
      // if overflow in nhi has occurred during carry over,
      // propagate that overflow to y and exit to increment z
      // otherwise return
      counter[1] += nhi;
      if(nhi != 0) {
        if (nhi <= counter[1]) {
          return;
        }
      }
    } else {
      // if overflow in y has occurred during addition,
      // exit to increment z
      // otherwise return
      counter[1] += nhi;
      if (nhi <= counter[1]) {
        return;
      }
    }
    if (++counter[2])
      return;
    ++counter[3];
  }

/**
   * Function that Skips one 128 bit number in a subsequence
   */
  C10_HOST_DEVICE inline void incr() {
    if (++counter[0])
      return;
    if (++counter[1])
      return;
    if (++counter[2]) {
      return;
    }
    ++counter[3];
  }

private:
  detail::UINT4 counter;
  detail::UINT4 output;
  detail::UINT2 key;
  uint32_t STATE;

C10_HOST_DEVICE inline uint32_t mulhilo32(uint32_t a, uint32_t b,
                                    uint32_t *result_high) {
    #ifdef __CUDA_ARCH__
      *result_high = __umulhi(a, b);
      return a*b;
    #else
      const uint64_t product = static_cast<uint64_t>(a) * b;
      *result_high = static_cast<uint32_t>(product >> 32);
      return static_cast<uint32_t>(product);
    #endif
  }

C10_HOST_DEVICE inline detail::UINT4 single_round(detail::UINT4 ctr, detail::UINT2 in_key) {
    uint32_t hi0;
    uint32_t hi1;
    uint32_t lo0 = mulhilo32(kPhiloxSA, ctr[0], &hi0);
    uint32_t lo1 = mulhilo32(kPhiloxSB, ctr[2], &hi1);
    detail::UINT4 ret;
    ret[0] = hi1 ^ ctr[1] ^ in_key[0];
    ret[1] = lo1;
    ret[2] = hi0 ^ ctr[3] ^ in_key[1];
    ret[3] = lo0;
    return ret;
  }
  static const uint32_t kPhilox10A = 0x9E3779B9;
  static const uint32_t kPhilox10B = 0xBB67AE85;
  static const uint32_t kPhiloxSA = 0xD2511F53;
  static const uint32_t kPhiloxSB = 0xCD9E8D57;
};

typedef philox_engine Philox4_32_10;

} // namespace at

cmd: