device/tensor_fill.h
Go to the documentation of this file.
/***************************************************************************************************
* Copyright (c) 2017-2019, NVIDIA CORPORATION. All rights reserved.
*
* Redistribution and use in source and binary forms, with or without modification, are permitted
* provided that the following conditions are met:
* * Redistributions of source code must retain the above copyright notice, this list of
* conditions and the following disclaimer.
* * Redistributions in binary form must reproduce the above copyright notice, this list of
* conditions and the following disclaimer in the documentation and/or other materials
* provided with the distribution.
* * Neither the name of the NVIDIA CORPORATION nor the names of its contributors may be used
* to endorse or promote products derived from this software without specific prior written
* permission.
*
* THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND ANY EXPRESS OR
* IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND
* FITNESS FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE
* FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING,
* BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS;
* OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT,
* STRICT LIABILITY, OR TOR (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
* OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
*
**************************************************************************************************/
/* \file
\brief Defines device-side elementwise operations on TensorView. Note, the operations defined
in this header are not specialized for any particular data layout and are therefore not
intended to offer the best possible performance. Rather, they are intended to be generic
reference implementations to support the CUTLASS unit tests.
*/
#pragma once
#if !defined(__CUDACC_RTC__)
// Standard Library includes
#include <utility>
#include <cstdlib>
#include <cmath>
#include <type_traits>
#include <cstdint>
#endif
// CUDA includes
#include <cublas_v2.h>
#include <curand_kernel.h>
// Cutlass includes
#include "cutlass/cutlass.h"
#include "cutlass/array.h"
#include "cutlass/tensor_view.h"
#include "cutlass/util/reference/device/tensor_foreach.h"
#include "cutlass/util/distribution.h"
namespace cutlass {
namespace reference {
namespace device {
namespace detail {
template <typename FloatType>
CUTLASS_DEVICE
FloatType random_normal_float(curandState_t *state) {
return curand_normal(state);
}
template <>
CUTLASS_DEVICE
double random_normal_float<double>(curandState_t *state) {
return curand_normal_double(state);
}
template <typename FloatType>
CUTLASS_DEVICE
FloatType random_uniform_float(curandState_t *state) {
return curand_uniform(state);
}
template <>
CUTLASS_DEVICE
double random_uniform_float<double>(curandState_t *state) {
return curand_uniform_double(state);
}
template <typename Element>
struct RandomGaussianFunc {
using FloatType = typename std::conditional<(sizeof(Element) > 4), double, float>::type;
using IntType = typename std::conditional<(sizeof(Element) > 4), int64_t, int>::type;
struct Params {
//
// Data members
//
uint64_t seed;
FloatType mean;
FloatType stddev;
int int_scale;
//
// Methods
//
Params(
uint64_t seed_ = 0,
Element mean_ = 0,
Element stddev_ = 1,
int int_scale_ = -1
):
seed(seed_),
mean(static_cast<FloatType>(mean_)),
stddev(static_cast<FloatType>(stddev_)),
int_scale(int_scale_) {
}
};
//
// Data members
//
Params params;
curandState_t rng_state;
//
// Methods
//
CUTLASS_DEVICE
RandomGaussianFunc(Params const &params): params(params) {
uint64_t gtid = threadIdx.x + blockIdx.x * blockDim.x;
curand_init(params.seed, gtid, 0, &rng_state);
}
CUTLASS_DEVICE
Element operator()() {
FloatType rnd = random_normal_float<FloatType>(&rng_state);
rnd = params.mean + params.stddev * rnd;
Element result;
if (params.int_scale >= 0) {
rnd = FloatType(IntType(rnd * FloatType(IntType(1) << params.int_scale)));
result = Element(rnd / FloatType(IntType(1) << params.int_scale));
}
else {
result = Element(rnd);
}
return result;
}
};
template <
typename Element,
typename Layout>
struct TensorFillRandomGaussianFunc {
using TensorView = TensorView<Element, Layout>;
typedef typename TensorView::Element T;
typedef typename TensorView::TensorCoord TensorCoord;
using RandomFunc = RandomGaussianFunc<Element>;
struct Params {
//
// Data members
//
TensorView view;
typename RandomFunc::Params random;
//
// Methods
//
Params(
TensorView view_ = TensorView(),
typename RandomFunc::Params random_ = typename RandomFunc::Params()
):
view(view_), random(random_) {
}
};
//
// Data members
//
Params params;
RandomFunc random;
//
// Methods
//
CUTLASS_DEVICE
TensorFillRandomGaussianFunc(Params const &params): params(params), random(params.random) {
}
CUTLASS_DEVICE
void operator()(TensorCoord const &coord) {
params.view.at(coord) = random();
}
};
} // namespace detail
template <
typename Element,
typename Layout>
void TensorFillRandomGaussian(
TensorView<Element, Layout> view,
uint64_t seed,
Element mean = Element(0),
Element stddev = Element(1),
int bits = -1) {
using RandomFunc = detail::RandomGaussianFunc<Element>;
using Func = detail::TensorFillRandomGaussianFunc<Element, Layout>;
using Params = typename Func::Params;
TensorForEach<Func, Layout::kRank, Params>(
view.extent(),
Params(view, typename RandomFunc::Params(seed, mean, stddev, bits))
);
}
template <typename Element>
void BlockFillRandomGaussian(
Element *ptr,
size_t capacity,
uint64_t seed,
Element mean = Element(0),
Element stddev = Element(1),
int bits = -1) {
using RandomFunc = detail::RandomGaussianFunc<Element>;
typename RandomFunc::Params params(seed, mean, stddev, bits);
BlockForEach<Element, RandomFunc>(ptr, capacity, params);
}
namespace detail {
template <typename Element>
struct RandomUniformFunc {
using FloatType = typename std::conditional<
(sizeof(Element) > 4),
double,
float>::type;
using IntType = typename std::conditional<
(sizeof(Element) > 4),
int64_t,
int>::type;
struct Params {
//
// Data members
//
uint64_t seed;
FloatType range;
FloatType min;
int int_scale;
CUTLASS_HOST_DEVICE
Params() { }
//
// Methods
//
Params(
uint64_t seed_ = 0,
Element max = 1,
Element min_ = 0,
int int_scale_ = -1
):
seed(seed_),
range(static_cast<FloatType>(max - min_)),
min(static_cast<FloatType>(min_)),
int_scale(int_scale_) {
}
};
//
// Data members
//
Params params;
curandState_t rng_state;
//
// Methods
//
CUTLASS_DEVICE
RandomUniformFunc(Params const &params): params(params) {
uint64_t gtid = threadIdx.x + blockIdx.x * blockDim.x;
curand_init(params.seed, gtid, 0, &rng_state);
}
CUTLASS_DEVICE
Element operator()() {
FloatType rnd = random_uniform_float<FloatType>(&rng_state);
rnd = params.min + params.range * rnd;
// Random values are cast to integer after scaling by a power of two to facilitate error
// testing
Element result;
if (params.int_scale >= 0) {
rnd = FloatType(IntType(rnd * FloatType(IntType(1) << params.int_scale)));
result = Element(rnd / FloatType(IntType(1) << params.int_scale));
}
else {
result = Element(rnd);
}
return result;
}
};
template <
typename Element,
typename Layout>
struct TensorFillRandomUniformFunc {
using TensorView = TensorView<Element, Layout>;
typedef typename TensorView::Element T;
typedef typename TensorView::TensorCoord TensorCoord;
using RandomFunc = RandomUniformFunc<Element>;
struct Params {
//
// Data members
//
TensorView view;
typename RandomFunc::Params random;
CUTLASS_HOST_DEVICE
Params() { }
//
// Methods
//
Params(
TensorView view_ = TensorView(),
typename RandomFunc::Params random_ = RandomFunc::Params()
):
view(view_), random(random_) {
}
};
//
// Data members
//
Params params;
RandomFunc random;
//
// Methods
//
CUTLASS_DEVICE
TensorFillRandomUniformFunc(Params const &params): params(params), random(params.random) {
}
CUTLASS_DEVICE
void operator()(TensorCoord const &coord) {
params.view.at(coord) = random();
}
};
} // namespace detail
template <
typename Element,
typename Layout>
void TensorFillRandomUniform(
TensorView<Element, Layout> view,
uint64_t seed,
Element max = Element(1),
Element min = Element(0),
int bits = -1) {
using RandomFunc = detail::RandomUniformFunc<Element>;
using Func = detail::TensorFillRandomUniformFunc<Element, Layout>;
using Params = typename Func::Params;
typename RandomFunc::Params random(seed, max, min, bits);
TensorForEach<Func, Layout::kRank, Params>(
view.size(),
Params(view, random)
);
}
template <typename Element>
void BlockFillRandomUniform(
Element *ptr,
size_t capacity,
uint64_t seed,
Element max = Element(1),
Element min = Element(0),
int bits = -1) {
using RandomFunc = detail::RandomUniformFunc<Element>;
typename RandomFunc::Params params(seed, max, min, bits);
BlockForEach<Element, RandomFunc>(ptr, capacity, params);
}
namespace detail {
template <
typename Element,
typename Layout>
struct TensorFillDiagonalFunc {
using TensorView = TensorView<Element, Layout>;
typedef typename TensorView::Element T;
typedef typename TensorView::TensorCoord TensorCoord;
struct Params {
//
// Data members
//
TensorView view;
Element diag;
Element other;
CUTLASS_HOST_DEVICE
Params() { }
//
// Methods
//
Params(
TensorView view_ = TensorView(),
Element diag_ = Element(1),
Element other_ = Element(0)
):
view(view_), diag(diag_), other(other_) {
}
};
//
// Data members
//
Params params;
//
// Methods
//
CUTLASS_DEVICE
TensorFillDiagonalFunc(Params const &params): params(params) {
}
CUTLASS_DEVICE
void operator()(TensorCoord const &coord) {
bool is_diag = true;
CUTLASS_PRAGMA_UNROLL
for (int i = 1; i < Layout::kRank; ++i) {
if (coord[i] != coord[i - 1]) {
is_diag = false;
break;
}
}
params.view.at(coord) = (is_diag ? params.diag : params.other);
}
};
} // namespace detail
template <
typename Element,
typename Layout>
void TensorFillDiagonal(
TensorView<Element, Layout> view,
Element diag = Element(1),
Element other = Element(0)) {
typedef detail::TensorFillDiagonalFunc<Element, Layout> Func;
typedef typename Func::Params Params;
TensorForEach<Func, Layout::kRank, Params>(
view.size(),
Params(view, diag, other)
);
}
template <
typename Element,
typename Layout>
void TensorFill(
TensorView<Element, Layout> view,
Element val = Element(0)) {
TensorFillDiagonal(view, val, val);
}
template <
typename Element,
typename Layout>
void TensorFillIdentity(
TensorView<Element, Layout> view) {
TensorFillDiagonal(view, Element(1), Element(0));
}
namespace detail {
template <
typename Element,
typename Layout>
struct TensorUpdateDiagonalFunc {
using TensorView = TensorView<Element, Layout>;
typedef typename TensorView::Element T;
typedef typename TensorView::TensorCoord TensorCoord;
struct Params {
//
// Data members
//
TensorView view;
Element diag;
CUTLASS_HOST_DEVICE
Params() { }
//
// Methods
//
Params(
TensorView view_ = TensorView(),
Element diag_ = Element(1)
):
view(view_), diag(diag_) {
}
};
//
// Data members
//
Params params;
//
// Methods
//
CUTLASS_DEVICE
TensorUpdateDiagonalFunc(Params const &params): params(params) {
}
CUTLASS_DEVICE
void operator()(TensorCoord const &coord) {
bool is_diag = true;
CUTLASS_PRAGMA_UNROLL
for (int i = 1; i < Layout::kRank; ++i) {
if (coord[i] != coord[i - 1]) {
is_diag = false;
break;
}
}
if (is_diag) {
params.view.at(coord) = params.diag;
}
}
};
} // namespace detail
template <
typename Element,
typename Layout>
void TensorUpdateDiagonal(
TensorView<Element, Layout> view,
Element diag = Element(1)) {
typedef detail::TensorUpdateDiagonalFunc<Element, Layout> Func;
typedef typename Func::Params Params;
TensorForEach<Func, Layout::kRank, Params>(
view.size(),
Params(view, diag)
);
}
namespace detail {
template <
typename Element,
typename Layout>
struct TensorUpdateOffDiagonalFunc {
using TensorView = TensorView<Element, Layout>;
typedef typename TensorView::Element T;
typedef typename TensorView::TensorCoord TensorCoord;
struct Params {
//
// Data members
//
TensorView view;
Element other;
CUTLASS_HOST_DEVICE
Params() { }
//
// Methods
//
Params(
TensorView view_ = TensorView(),
Element other_ = Element(0)
):
view(view_), other(other_) {
}
};
//
// Data members
//
Params params;
//
// Methods
//
CUTLASS_DEVICE
TensorUpdateOffDiagonalFunc(Params const &params): params(params) {
}
CUTLASS_DEVICE
void operator()(TensorCoord const &coord) {
bool is_diag = true;
CUTLASS_PRAGMA_UNROLL
for (int i = 1; i < Layout::kRank; ++i) {
if (coord[i] != coord[i - 1]) {
is_diag = false;
break;
}
}
if (!is_diag) {
params.view.at(coord) = params.other;
}
}
};
} // namespace detail
template <
typename Element,
typename Layout>
void TensorUpdateOffDiagonal(
TensorView<Element, Layout> view,
Element other = Element(1)) {
typedef detail::TensorUpdateOffDiagonalFunc<Element, Layout> Func;
typedef typename Func::Params Params;
TensorForEach<Func, Layout::kRank, Params>(
view.size(),
Params(view, other)
);
}
namespace detail {
template <
typename Element,
typename Layout>
struct TensorFillLinearFunc {
using TensorView = TensorView<Element, Layout>;
typedef typename TensorView::Element T;
typedef typename TensorView::TensorCoord TensorCoord;
struct Params {
//
// Data members
//
TensorView view;
Array<Element, Layout::kRank> v;
Element s;
CUTLASS_HOST_DEVICE
Params() { }
//
// Methods
//
Params(
TensorView view_,
Array<Element, Layout::kRank> const & v_,
Element s_ = Element(0)
):
view(view_), v(v_), s(s_) {
}
};
//
// Data members
//
Params params;
//
// Methods
//
CUTLASS_DEVICE
TensorFillLinearFunc(Params const &params): params(params) {
}
CUTLASS_DEVICE
void operator()(TensorCoord const &coord) {
Element sum = params.s;
CUTLASS_PRAGMA_UNROLL
for (int i = 0; i < Layout::kRank; ++i) {
sum += params.v[i] * Element(coord[i]);
}
params.view.at(coord) = sum;
}
};
} // namespace detail
template <
typename Element,
typename Layout>
void TensorFillLinear(
TensorView<Element, Layout> view,
Array<Element, Layout::kRank> const & v,
Element s = Element(0)) {
using Func = detail::TensorFillLinearFunc<Element, Layout>;
using Params = typename Func::Params;
TensorForEach<Func, Layout::kRank, Params>(
view.size(),
Params(view, v, s)
);
}
template <
typename Element
>
void BlockFillSequential(
Element *ptr,
int64_t capacity,
Element v = Element(1),
Element s = Element(0)) {
}
template <
typename Element
>
void BlockFillRandom(
Element *ptr,
size_t capacity,
uint64_t seed,
Distribution dist) {
if (dist.kind == Distribution::Gaussian) {
BlockFillRandomGaussian<Element>(
ptr,
capacity,
seed,
static_cast<Element>(dist.gaussian.mean),
static_cast<Element>(dist.gaussian.stddev),
dist.int_scale);
}
else if (dist.kind == Distribution::Uniform) {
BlockFillRandomUniform<Element>(
ptr,
capacity,
seed,
static_cast<Element>(dist.uniform.max),
static_cast<Element>(dist.uniform.min),
dist.int_scale);
}
}
namespace detail {
template <
typename Element,
typename Layout>
struct TensorCopyDiagonalInFunc {
using TensorView = TensorView<Element, Layout>;
typedef typename TensorView::Element T;
typedef typename TensorView::TensorCoord TensorCoord;
struct Params {
//
// Data members
//
TensorView view;
Element const *ptr;
CUTLASS_HOST_DEVICE
Params() { }
//
// Methods
//
Params(
TensorView view_,
Element const *ptr_
):
view(view_), ptr(ptr_) {
}
};
//
// Data members
//
Params params;
//
// Methods
//
CUTLASS_DEVICE
TensorCopyDiagonalInFunc(Params const &params): params(params) {
}
CUTLASS_DEVICE
void operator()(TensorCoord const &coord) {
bool is_diagonal = true;
CUTLASS_PRAGMA_UNROLL
for (int i = 1; i < Layout::kRank; ++i) {
if (coord[i] != coord[0]) {
is_diagonal = false;
}
}
if (is_diagonal) {
params.view.at(coord) = params.ptr[coord[0]];
}
}
};
} // namespace detail
template <
typename Element,
typename Layout>
void TensorCopyDiagonalIn(
TensorView<Element, Layout> view,
Element const *ptr) {
using Func = detail::TensorCopyDiagonalInFunc<Element, Layout>;
using Params = typename Func::Params;
TensorForEach<Func, Layout::kRank, Params>(
view.size(),
Params(view, ptr)
);
}
namespace detail {
template <
typename Element,
typename Layout>
struct TensorCopyDiagonalOutFunc {
using TensorView = TensorView<Element, Layout>;
typedef typename TensorView::Element T;
typedef typename TensorView::TensorCoord TensorCoord;
struct Params {
//
// Data members
//
TensorView view;
Element *ptr;
CUTLASS_HOST_DEVICE
Params() { }
//
// Methods
//
Params(
TensorView view_,
Element *ptr_
):
view(view_), ptr(ptr_) {
}
};
//
// Data members
//
Params params;
//
// Methods
//
CUTLASS_DEVICE
TensorCopyDiagonalOutFunc(Params const &params): params(params) {
}
CUTLASS_DEVICE
void operator()(TensorCoord const &coord) {
bool is_diagonal = true;
CUTLASS_PRAGMA_UNROLL
for (int i = 1; i < Layout::kRank; ++i) {
if (coord[i] != coord[0]) {
is_diagonal = false;
}
}
if (is_diagonal) {
params.ptr[coord[0]] = params.view.at(coord);
}
}
};
} // namespace detail
template <
typename Element,
typename Layout>
void TensorCopyDiagonalOut(
Element *ptr,
TensorView<Element, Layout> view) {
using Func = detail::TensorCopyDiagonalOutFunc<Element, Layout>;
using Params = typename Func::Params;
TensorForEach<Func, Layout::kRank, Params>(
view.size(),
Params(view, ptr)
);
}
} // namespace device
} // namespace reference
} // namespace cutlass
cutlass::reference::device::detail::TensorFillLinearFunc::TensorCoord
TensorView::TensorCoord TensorCoord
Coordinate in tensor&#39;s index space.
Definition: device/tensor_fill.h:866
cutlass::reference::device::detail::TensorUpdateOffDiagonalFunc
Computes a random Gaussian distribution.
Definition: device/tensor_fill.h:751
cutlass::reference::device::detail::TensorFillLinearFunc::Params::Params
CUTLASS_HOST_DEVICE Params()
Default ctor.
Definition: device/tensor_fill.h:881
cutlass::reference::device::detail::TensorUpdateDiagonalFunc::Params::Params
Params(TensorView view_=TensorView(), Element diag_=Element(1))
Construction of Gaussian RNG functor.
Definition: device/tensor_fill.h:675
cutlass::platform::max
CUTLASS_HOST_DEVICE constexpr const T & max(const T &a, const T &b)
std::max
Definition: platform.h:189
cutlass::reference::device::detail::RandomGaussianFunc::RandomGaussianFunc
CUTLASS_DEVICE RandomGaussianFunc(Params const &params)
Device-side initialization of RNG.
Definition: device/tensor_fill.h:145
cutlass::reference::device::detail::TensorFillRandomUniformFunc::TensorFillRandomUniformFunc
CUTLASS_DEVICE TensorFillRandomUniformFunc(Params const &params)
Device-side initialization of RNG.
Definition: device/tensor_fill.h:444
cutlass::reference::device::detail::TensorCopyDiagonalInFunc::Params::Params
CUTLASS_HOST_DEVICE Params()
Default ctor.
Definition: device/tensor_fill.h:1031
cutlass::reference::device::detail::TensorUpdateOffDiagonalFunc::TensorCoord
TensorView::TensorCoord TensorCoord
Coordinate in tensor&#39;s index space.
Definition: device/tensor_fill.h:760
cutlass
Definition: aligned_buffer.h:35
cutlass::Distribution::Uniform
Definition: distribution.h:40
cutlass::reference::device::detail::TensorFillRandomUniformFunc::random
RandomFunc random
Definition: device/tensor_fill.h:436
cutlass::reference::device::TensorCopyDiagonalOut
void TensorCopyDiagonalOut(Element *ptr, TensorView< Element, Layout > view)
Copies the diagonal of a tensor into a dense buffer in host memory.
Definition: device/tensor_fill.h:1193
cutlass::reference::device::detail::RandomUniformFunc::IntType
typename std::conditional< (sizeof(Element) > 4), int64_t, int >::type IntType
Definition: device/tensor_fill.h:303
cutlass::reference::device::detail::RandomUniformFunc::operator()
CUTLASS_DEVICE Element operator()()
Compute random value and update RNG state.
Definition: device/tensor_fill.h:365
cutlass::reference::device::detail::TensorFillRandomUniformFunc::Params::Params
CUTLASS_HOST_DEVICE Params()
Default ctor.
Definition: device/tensor_fill.h:415
cutlass::Distribution::Gaussian
Definition: distribution.h:40
cutlass::reference::device::detail::TensorUpdateDiagonalFunc::Params::view
TensorView view
Definition: device/tensor_fill.h:663
cutlass::Distribution::uniform
struct cutlass::Distribution::@18::@20 uniform
Uniform distribution.
cutlass::reference::device::detail::TensorCopyDiagonalInFunc::Params
Parameters structure.
Definition: device/tensor_fill.h:1020
cutlass::reference::device::detail::TensorFillLinearFunc::T
TensorView::Element T
Scalar type.
Definition: device/tensor_fill.h:863
cutlass::reference::device::detail::TensorCopyDiagonalInFunc::Params::ptr
Element const * ptr
Definition: device/tensor_fill.h:1027
cutlass::reference::device::detail::TensorFillRandomGaussianFunc::Params::view
TensorView view
Definition: device/tensor_fill.h:196
cutlass::reference::device::detail::TensorCopyDiagonalInFunc::TensorCoord
TensorView::TensorCoord TensorCoord
Coordinate in tensor&#39;s index space.
Definition: device/tensor_fill.h:1017
cutlass::reference::device::detail::RandomGaussianFunc::Params
Parameters structure.
Definition: device/tensor_fill.h:99
cutlass::Distribution::kind
Kind kind
Active variant kind.
Definition: distribution.h:64
cutlass::reference::device::detail::TensorFillRandomUniformFunc::Params::Params
Params(TensorView view_=TensorView(), typename RandomFunc::Params random_=RandomFunc::Params())
Construction of Gaussian RNG functor.
Definition: device/tensor_fill.h:422
cutlass::reference::device::TensorFillIdentity
void TensorFillIdentity(TensorView< Element, Layout > view)
Fills a tensor&#39;s diagonal with 1 and 0 everywhere else.
Definition: device/tensor_fill.h:630
cutlass::TensorView::extent
CUTLASS_HOST_DEVICE TensorCoord const & extent() const
Returns the extent of the view (the size along each logical dimension).
Definition: tensor_view.h:167
cutlass::reference::device::detail::TensorUpdateDiagonalFunc
Computes a random Gaussian distribution.
Definition: device/tensor_fill.h:645
cutlass::reference::device::detail::RandomUniformFunc::Params::int_scale
int int_scale
Definition: device/tensor_fill.h:315
cutlass::reference::device::detail::TensorFillRandomUniformFunc::T
TensorView::Element T
Scalar type.
Definition: device/tensor_fill.h:396
cutlass::reference::device::detail::TensorCopyDiagonalOutFunc::Params::Params
Params(TensorView view_, Element *ptr_)
Construction of Gaussian RNG functor.
Definition: device/tensor_fill.h:1142
cutlass::reference::device::detail::TensorCopyDiagonalInFunc::params
Params params
Parameters object.
Definition: device/tensor_fill.h:1052
cutlass::reference::device::detail::TensorFillRandomUniformFunc::Params::random
RandomFunc::Params random
Definition: device/tensor_fill.h:411
cutlass::Distribution::gaussian
struct cutlass::Distribution::@18::@21 gaussian
Gaussian distribution.
cutlass::reference::device::detail::RandomUniformFunc::Params::min
FloatType min
Definition: device/tensor_fill.h:314
cutlass::reference::device::detail::RandomGaussianFunc
Definition: device/tensor_fill.h:93
cutlass::reference::device::detail::TensorUpdateDiagonalFunc::T
TensorView::Element T
Scalar type.
Definition: device/tensor_fill.h:651
cutlass::reference::device::detail::TensorUpdateDiagonalFunc::TensorUpdateDiagonalFunc
CUTLASS_DEVICE TensorUpdateDiagonalFunc(Params const &params)
Device-side initialization of RNG.
Definition: device/tensor_fill.h:697
cutlass::reference::device::detail::TensorFillLinearFunc::TensorFillLinearFunc
CUTLASS_DEVICE TensorFillLinearFunc(Params const &params)
Device-side initialization of RNG.
Definition: device/tensor_fill.h:911
cutlass::reference::device::TensorCopyDiagonalIn
void TensorCopyDiagonalIn(TensorView< Element, Layout > view, Element const *ptr)
Copies a diagonal in from host memory without modifying off-diagonal elements.
Definition: device/tensor_fill.h:1089
cutlass::reference::device::detail::RandomGaussianFunc::rng_state
curandState_t rng_state
RNG state object.
Definition: device/tensor_fill.h:137
cutlass::reference::device::detail::TensorFillLinearFunc::operator()
CUTLASS_DEVICE void operator()(TensorCoord const &coord)
Compute random value and update RNG state.
Definition: device/tensor_fill.h:917
tensor_foreach.h
cutlass::reference::device::detail::RandomUniformFunc::Params::Params
Params(uint64_t seed_=0, Element max=1, Element min_=0, int int_scale_=-1)
Construction of Gaussian RNG functor.
Definition: device/tensor_fill.h:326
tensor_view.h
Defines a structure containing strides and a pointer to tensor data.
cutlass::reference::device::detail::TensorFillRandomGaussianFunc::random
RandomFunc random
Definition: device/tensor_fill.h:218
cutlass::reference::device::detail::RandomGaussianFunc::Params::seed
uint64_t seed
Definition: device/tensor_fill.h:105
cutlass::reference::device::detail::random_normal_float< double >
CUTLASS_DEVICE double random_normal_float< double >(curandState_t *state)
Definition: device/tensor_fill.h:76
cutlass::FloatType
Defines a floating-point type based on the number of exponent and mantissa bits.
Definition: numeric_types.h:144
cutlass::reference::device::detail::RandomGaussianFunc::FloatType
typename std::conditional<(sizeof(Element) > 4), double, float >::type FloatType
Definition: device/tensor_fill.h:95
cutlass::reference::device::detail::TensorFillLinearFunc::Params::view
TensorView view
Definition: device/tensor_fill.h:875
cutlass::TensorView< Element, Layout >::Element
Element Element
Data type of individual access.
Definition: tensor_view.h:72
cutlass::reference::device::detail::TensorCopyDiagonalInFunc::Params::Params
Params(TensorView view_, Element const *ptr_)
Construction of Gaussian RNG functor.
Definition: device/tensor_fill.h:1038
cutlass::reference::device::detail::TensorFillDiagonalFunc::Params::Params
Params(TensorView view_=TensorView(), Element diag_=Element(1), Element other_=Element(0))
Construction of Gaussian RNG functor.
Definition: device/tensor_fill.h:544
cutlass::reference::device::detail::RandomUniformFunc::Params::seed
uint64_t seed
Definition: device/tensor_fill.h:312
cutlass::reference::device::BlockFillSequential
void BlockFillSequential(Element *ptr, int64_t capacity, Element v=Element(1), Element s=Element(0))
Fills a block of data with sequential elements.
Definition: device/tensor_fill.h:958
array.h
Statically sized array of elements that accommodates all CUTLASS-supported numeric types and is safe ...
CUTLASS_PRAGMA_UNROLL
#define CUTLASS_PRAGMA_UNROLL
Definition: cutlass.h:110
cutlass::reference::device::detail::RandomUniformFunc::Params
Parameters structure.
Definition: device/tensor_fill.h:306
cutlass::reference::device::detail::TensorFillLinearFunc
Computes a random Gaussian distribution.
Definition: device/tensor_fill.h:857
cutlass::reference::device::detail::RandomGaussianFunc::Params::int_scale
int int_scale
Definition: device/tensor_fill.h:108
cutlass::reference::device::TensorFillRandomGaussian
void TensorFillRandomGaussian(TensorView< Element, Layout > view, uint64_t seed, Element mean=Element(0), Element stddev=Element(1), int bits=-1)
Fills a tensor with random values with a Gaussian distribution.
Definition: device/tensor_fill.h:246
cutlass::reference::device::detail::TensorFillDiagonalFunc::TensorFillDiagonalFunc
CUTLASS_DEVICE TensorFillDiagonalFunc(Params const &params)
Device-side initialization of RNG.
Definition: device/tensor_fill.h:567
cutlass::reference::device::detail::TensorFillRandomGaussianFunc::params
Params params
Definition: device/tensor_fill.h:217
cutlass::reference::device::BlockFillRandomUniform
void BlockFillRandomUniform(Element *ptr, size_t capacity, uint64_t seed, Element max=Element(1), Element min=Element(0), int bits=-1)
Fills a tensor with random values with a uniform random distribution.
Definition: device/tensor_fill.h:488
cutlass::reference::device::detail::TensorFillRandomGaussianFunc::Params
Parameters structure.
Definition: device/tensor_fill.h:190
cutlass::reference::device::detail::TensorCopyDiagonalInFunc::Params::view
TensorView view
Definition: device/tensor_fill.h:1026
cutlass::reference::device::detail::TensorFillLinearFunc::Params
Parameters structure.
Definition: device/tensor_fill.h:869
cutlass::reference::device::detail::TensorFillLinearFunc::Params::Params
Params(TensorView view_, Array< Element, Layout::kRank > const &v_, Element s_=Element(0))
Construction of Gaussian RNG functor.
Definition: device/tensor_fill.h:888
cutlass::TensorView< Element, Layout >
cutlass::reference::device::detail::TensorFillRandomUniformFunc::params
Params params
Definition: device/tensor_fill.h:435
cutlass::reference::device::TensorFillDiagonal
void TensorFillDiagonal(TensorView< Element, Layout > view, Element diag=Element(1), Element other=Element(0))
Fills a tensor everywhere with a unique value for its diagonal.
Definition: device/tensor_fill.h:597
cutlass::TensorView< Element, Layout >::TensorCoord
typename Layout::TensorCoord TensorCoord
Coordinate in logical tensor space.
Definition: tensor_view.h:87
cutlass::reference::device::detail::TensorFillLinearFunc::Params::s
Element s
Definition: device/tensor_fill.h:877
cutlass::reference::device::detail::TensorCopyDiagonalOutFunc::T
TensorView::Element T
Scalar type.
Definition: device/tensor_fill.h:1118
cutlass::reference::device::detail::TensorUpdateOffDiagonalFunc::params
Params params
Parameters object.
Definition: device/tensor_fill.h:795
cutlass::reference::device::detail::TensorFillDiagonalFunc
Computes a random Gaussian distribution.
Definition: device/tensor_fill.h:513
cutlass::reference::device::detail::TensorFillDiagonalFunc::Params
Parameters structure.
Definition: device/tensor_fill.h:525
cutlass::reference::device::detail::TensorCopyDiagonalInFunc
Computes a random Gaussian distribution.
Definition: device/tensor_fill.h:1008
cutlass::reference::device::detail::RandomGaussianFunc::Params::mean
FloatType mean
Definition: device/tensor_fill.h:106
cutlass::reference::device::detail::TensorUpdateDiagonalFunc::Params::diag
Element diag
Definition: device/tensor_fill.h:664
cutlass::reference::device::detail::TensorCopyDiagonalInFunc::TensorCopyDiagonalInFunc
CUTLASS_DEVICE TensorCopyDiagonalInFunc(Params const &params)
Device-side initialization of RNG.
Definition: device/tensor_fill.h:1060
cutlass::reference::device::detail::TensorUpdateDiagonalFunc::TensorCoord
TensorView::TensorCoord TensorCoord
Coordinate in tensor&#39;s index space.
Definition: device/tensor_fill.h:654
cutlass::reference::device::detail::TensorFillDiagonalFunc::operator()
CUTLASS_DEVICE void operator()(TensorCoord const &coord)
Compute random value and update RNG state.
Definition: device/tensor_fill.h:573
cutlass::reference::device::TensorFill
void TensorFill(TensorView< Element, Layout > view, Element val=Element(0))
Fills a tensor with a uniform value.
Definition: device/tensor_fill.h:617
distribution.h
This header contains a class to parametrize a statistical distribution function.
cutlass::reference::device::detail::RandomUniformFunc
Computes a random Gaussian distribution.
Definition: device/tensor_fill.h:293
cutlass::reference::device::detail::RandomGaussianFunc::params
Params params
Parameters object.
Definition: device/tensor_fill.h:134
cutlass::reference::device::detail::TensorUpdateDiagonalFunc::operator()
CUTLASS_DEVICE void operator()(TensorCoord const &coord)
Compute random value and update RNG state.
Definition: device/tensor_fill.h:703
cutlass::reference::device::detail::TensorFillRandomUniformFunc::operator()
CUTLASS_DEVICE void operator()(TensorCoord const &coord)
Compute random value and update RNG state.
Definition: device/tensor_fill.h:449
cutlass::reference::device::BlockFillRandomGaussian
void BlockFillRandomGaussian(Element *ptr, size_t capacity, uint64_t seed, Element mean=Element(0), Element stddev=Element(1), int bits=-1)
Fills a tensor with random values with a Gaussian distribution.
Definition: device/tensor_fill.h:269
cutlass::reference::device::detail::TensorCopyDiagonalOutFunc::TensorCoord
TensorView::TensorCoord TensorCoord
Coordinate in tensor&#39;s index space.
Definition: device/tensor_fill.h:1121
cutlass::reference::device::detail::TensorCopyDiagonalInFunc::T
TensorView::Element T
Scalar type.
Definition: device/tensor_fill.h:1014
cutlass::reference::device::detail::TensorFillRandomGaussianFunc::TensorCoord
TensorView::TensorCoord TensorCoord
Coordinate in tensor&#39;s index space.
Definition: device/tensor_fill.h:185
CUTLASS_HOST_DEVICE
#define CUTLASS_HOST_DEVICE
Definition: cutlass.h:89
cutlass::reference::device::detail::TensorUpdateOffDiagonalFunc::Params::other
Element other
Definition: device/tensor_fill.h:770
cutlass::reference::device::detail::TensorFillRandomGaussianFunc::operator()
CUTLASS_DEVICE void operator()(TensorCoord const &coord)
Compute random value and update RNG state.
Definition: device/tensor_fill.h:232
cutlass::reference::device::detail::TensorCopyDiagonalInFunc::operator()
CUTLASS_DEVICE void operator()(TensorCoord const &coord)
Only update the diagonal element.
Definition: device/tensor_fill.h:1066
cutlass::reference::device::detail::TensorCopyDiagonalOutFunc
Computes a random Gaussian distribution.
Definition: device/tensor_fill.h:1112
cutlass::platform::min
CUTLASS_HOST_DEVICE constexpr const T & min(const T &a, const T &b)
std::min
Definition: platform.h:183
cutlass::reference::device::detail::TensorCopyDiagonalOutFunc::Params::view
TensorView view
Definition: device/tensor_fill.h:1130
cutlass::reference::device::detail::TensorFillDiagonalFunc::Params::other
Element other
Definition: device/tensor_fill.h:533
cutlass::reference::device::detail::TensorUpdateOffDiagonalFunc::Params
Parameters structure.
Definition: device/tensor_fill.h:763
cutlass::reference::device::detail::RandomUniformFunc::FloatType
typename std::conditional< (sizeof(Element) > 4), double, float >::type FloatType
Definition: device/tensor_fill.h:298
cutlass::reference::device::TensorForEach
Launches a kernel calling a functor for each element in a tensor&#39;s index space.
Definition: device/tensor_foreach.h:39
cutlass::reference::device::detail::TensorUpdateDiagonalFunc::Params
Parameters structure.
Definition: device/tensor_fill.h:657
cutlass::reference::device::detail::TensorFillLinearFunc::Params::v
Array< Element, Layout::kRank > v
Definition: device/tensor_fill.h:876
cutlass::reference::device::TensorUpdateDiagonal
void TensorUpdateDiagonal(TensorView< Element, Layout > view, Element diag=Element(1))
Writes a uniform value to the diagonal of a tensor without modifying off-diagonal elements...
Definition: device/tensor_fill.h:729
cutlass::reference::device::detail::TensorFillRandomUniformFunc
Computes a random Gaussian distribution.
Definition: device/tensor_fill.h:390
cutlass::reference::device::detail::random_uniform_float< double >
CUTLASS_DEVICE double random_uniform_float< double >(curandState_t *state)
Definition: device/tensor_fill.h:88
cutlass::reference::device::detail::TensorUpdateOffDiagonalFunc::Params::view
TensorView view
Definition: device/tensor_fill.h:769
cutlass::reference::device::detail::random_normal_float
CUTLASS_DEVICE FloatType random_normal_float(curandState_t *state)
Definition: device/tensor_fill.h:70
cutlass::reference::device::detail::TensorUpdateOffDiagonalFunc::TensorUpdateOffDiagonalFunc
CUTLASS_DEVICE TensorUpdateOffDiagonalFunc(Params const &params)
Device-side initialization of RNG.
Definition: device/tensor_fill.h:803
cutlass::reference::device::detail::RandomUniformFunc::Params::range
FloatType range
Definition: device/tensor_fill.h:313
cutlass::reference::device::BlockFillRandom
void BlockFillRandom(Element *ptr, size_t capacity, uint64_t seed, Distribution dist)
Fills a block of data with sequential elements.
Definition: device/tensor_fill.h:973
cutlass::reference::device::detail::RandomGaussianFunc::Params::Params
Params(uint64_t seed_=0, Element mean_=0, Element stddev_=1, int int_scale_=-1)
Construction of Gaussian RNG functor.
Definition: device/tensor_fill.h:115
cutlass::reference::device::detail::TensorCopyDiagonalOutFunc::TensorCopyDiagonalOutFunc
CUTLASS_DEVICE TensorCopyDiagonalOutFunc(Params const &params)
Device-side initialization of RNG.
Definition: device/tensor_fill.h:1164
cutlass::reference::device::TensorFillLinear
void TensorFillLinear(TensorView< Element, Layout > view, Array< Element, Layout::kRank > const &v, Element s=Element(0))
Fills tensor with a linear combination of its coordinate and another vector.
Definition: device/tensor_fill.h:937
cutlass::reference::device::detail::RandomUniformFunc::RandomUniformFunc
CUTLASS_DEVICE RandomUniformFunc(Params const &params)
Device-side initialization of RNG.
Definition: device/tensor_fill.h:356
cutlass::reference::device::detail::TensorCopyDiagonalOutFunc::Params::ptr
Element * ptr
Definition: device/tensor_fill.h:1131
cutlass::reference::device::detail::TensorFillDiagonalFunc::params
Params params
Parameters object.
Definition: device/tensor_fill.h:559
cutlass::reference::device::detail::TensorUpdateDiagonalFunc::params
Params params
Parameters object.
Definition: device/tensor_fill.h:689
cutlass::reference::device::detail::TensorCopyDiagonalOutFunc::Params
Parameters structure.
Definition: device/tensor_fill.h:1124
cutlass::reference::device::TensorUpdateOffDiagonal
void TensorUpdateOffDiagonal(TensorView< Element, Layout > view, Element other=Element(1))
Writes a uniform value to all elements in the tensor without modifying diagonal elements.
Definition: device/tensor_fill.h:835
cutlass::reference::device::detail::RandomUniformFunc::params
Params params
Parameters object.
Definition: device/tensor_fill.h:345
cutlass::reference::device::detail::TensorUpdateOffDiagonalFunc::T
TensorView::Element T
Scalar type.
Definition: device/tensor_fill.h:757
cutlass::reference::device::detail::TensorFillRandomUniformFunc::TensorCoord
TensorView::TensorCoord TensorCoord
Coordinate in tensor&#39;s index space.
Definition: device/tensor_fill.h:399
cutlass::reference::device::detail::TensorFillDiagonalFunc::Params::diag
Element diag
Definition: device/tensor_fill.h:532
cutlass::reference::device::detail::RandomGaussianFunc::IntType
typename std::conditional<(sizeof(Element) > 4), int64_t, int >::type IntType
Definition: device/tensor_fill.h:96
cutlass::reference::device::detail::RandomUniformFunc::Params::Params
CUTLASS_HOST_DEVICE Params()
Default ctor.
Definition: device/tensor_fill.h:319
cutlass::TensorRef::at
CUTLASS_HOST_DEVICE Reference at(TensorCoord const &coord) const
Returns a reference to the element at a given Coord.
Definition: tensor_ref.h:307
cutlass::reference::device::detail::TensorCopyDiagonalOutFunc::Params::Params
CUTLASS_HOST_DEVICE Params()
Default ctor.
Definition: device/tensor_fill.h:1135
cutlass::reference::device::detail::TensorFillDiagonalFunc::Params::view
TensorView view
Definition: device/tensor_fill.h:531
cutlass::reference::device::detail::random_uniform_float
CUTLASS_DEVICE FloatType random_uniform_float(curandState_t *state)
Definition: device/tensor_fill.h:82
cutlass::reference::device::detail::TensorUpdateDiagonalFunc::Params::Params
CUTLASS_HOST_DEVICE Params()
Default ctor.
Definition: device/tensor_fill.h:668
cutlass::reference::device::TensorFillRandomUniform
void TensorFillRandomUniform(TensorView< Element, Layout > view, uint64_t seed, Element max=Element(1), Element min=Element(0), int bits=-1)
Fills a tensor with random values with a uniform random distribution.
Definition: device/tensor_fill.h:463
cutlass::reference::device::detail::TensorCopyDiagonalOutFunc::params
Params params
Parameters object.
Definition: device/tensor_fill.h:1156
cutlass::reference::device::detail::TensorFillDiagonalFunc::TensorCoord
TensorView::TensorCoord TensorCoord
Coordinate in tensor&#39;s index space.
Definition: device/tensor_fill.h:522
cutlass::reference::device::detail::TensorFillRandomUniformFunc::Params::view
TensorView view
Definition: device/tensor_fill.h:410
cutlass::reference::device::detail::RandomGaussianFunc::operator()
CUTLASS_DEVICE Element operator()()
Compute random value and update RNG state.
Definition: device/tensor_fill.h:154
cutlass::reference::device::BlockForEach
Definition: device/tensor_foreach.h:92
cutlass::reference::device::detail::TensorCopyDiagonalOutFunc::operator()
CUTLASS_DEVICE void operator()(TensorCoord const &coord)
Compute random value and update RNG state.
Definition: device/tensor_fill.h:1170
cutlass::Distribution
Distribution type.
Definition: distribution.h:38
cutlass::reference::device::detail::RandomUniformFunc::rng_state
curandState_t rng_state
RNG state object.
Definition: device/tensor_fill.h:348
cutlass::reference::device::detail::TensorFillRandomGaussianFunc
Computes a random Gaussian distribution.
Definition: device/tensor_fill.h:176
cutlass::reference::device::detail::TensorUpdateOffDiagonalFunc::Params::Params
CUTLASS_HOST_DEVICE Params()
Default ctor.
Definition: device/tensor_fill.h:774
cutlass::reference::device::detail::TensorFillDiagonalFunc::Params::Params
CUTLASS_HOST_DEVICE Params()
Default ctor.
Definition: device/tensor_fill.h:537
cutlass::Distribution::int_scale
int int_scale
Random values are cast to integer after scaling by this power of two.
Definition: distribution.h:67
cutlass::reference::device::detail::TensorFillDiagonalFunc::T
TensorView::Element T
Scalar type.
Definition: device/tensor_fill.h:519
cutlass::reference::device::detail::TensorFillRandomGaussianFunc::T
TensorView::Element T
Scalar type.
Definition: device/tensor_fill.h:182
cutlass.h
Basic include for CUTLASS.
cutlass::reference::device::detail::TensorFillRandomGaussianFunc::Params::Params
Params(TensorView view_=TensorView(), typename RandomFunc::Params random_=typename RandomFunc::Params())
Construction of Gaussian RNG functor.
Definition: device/tensor_fill.h:204
cutlass::reference::device::detail::TensorFillRandomGaussianFunc::TensorFillRandomGaussianFunc
CUTLASS_DEVICE TensorFillRandomGaussianFunc(Params const &params)
Device-side initialization of RNG.
Definition: device/tensor_fill.h:226
cutlass::reference::device::detail::TensorUpdateOffDiagonalFunc::operator()
CUTLASS_DEVICE void operator()(TensorCoord const &coord)
Compute random value and update RNG state.
Definition: device/tensor_fill.h:809
cutlass::reference::device::detail::TensorFillLinearFunc::params
Params params
Parameters object.
Definition: device/tensor_fill.h:903
cutlass::reference::device::detail::TensorUpdateOffDiagonalFunc::Params::Params
Params(TensorView view_=TensorView(), Element other_=Element(0))
Construction of Gaussian RNG functor.
Definition: device/tensor_fill.h:781
cutlass::reference::device::detail::TensorFillRandomUniformFunc::Params
Parameters structure.
Definition: device/tensor_fill.h:404
cutlass::reference::device::detail::TensorFillRandomGaussianFunc::Params::random
RandomFunc::Params random
Definition: device/tensor_fill.h:197
cutlass::reference::device::detail::RandomGaussianFunc::Params::stddev
FloatType stddev
Definition: device/tensor_fill.h:107