conversion_op.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.
*
**************************************************************************************************/
#pragma once
#include "cutlass/cutlass.h"
#include "cutlass/numeric_types.h"
#include "cutlass/array.h"
#include "cutlass/functional.h"
#include "cutlass/numeric_conversion.h"
namespace cutlass {
namespace epilogue {
namespace thread {
template <
typename ElementOutput_,
int Count,
typename ElementAccumulator_ = ElementOutput_,
FloatRoundStyle Round = FloatRoundStyle::round_to_nearest
>
class Convert {
public:
using ElementOutput = ElementOutput_;
using ElementAccumulator = ElementAccumulator_;
using ElementCompute = ElementAccumulator_;
static int const kCount = Count;
using FragmentOutput = Array<ElementOutput, kCount>;
using FragmentAccumulator = Array<ElementAccumulator, kCount>;
using ComputeFragment = FragmentAccumulator;
static FloatRoundStyle const kRound = Round;
struct Params {
//
// Methods
//
CUTLASS_HOST_DEVICE
Params() {}
};
public:
CUTLASS_HOST_DEVICE
Convert(Params const &params = Params()) {
}
CUTLASS_HOST_DEVICE
constexpr bool is_source_needed() const {
return false;
}
CUTLASS_HOST_DEVICE
constexpr bool is_source_ever_needed() const {
return false;
}
CUTLASS_HOST_DEVICE
FragmentOutput operator()(
FragmentAccumulator const &accumulator,
FragmentOutput const &source,
ElementCompute uniform = ElementCompute(0)) const {
// Convert to destination numeric type
NumericArrayConverter<ElementOutput, ElementAccumulator, kCount, Round> destination_converter;
return destination_converter(accumulator);
}
};
} // namespace thread
} // namespace epilogue
} // namespace cutlass
cutlass::epilogue::thread::Convert
Definition: conversion_op.h:53
cutlass::epilogue::thread::Convert::Convert
CUTLASS_HOST_DEVICE Convert(Params const &params=Params())
Constructs the function object, possibly loading from pointers in host memory.
Definition: conversion_op.h:83
cutlass::epilogue::thread::Convert::ComputeFragment
FragmentAccumulator ComputeFragment
Definition: conversion_op.h:64
cutlass
Definition: aligned_buffer.h:35
constexpr
#define constexpr
Definition: platform.h:137
cutlass::epilogue::thread::Convert::is_source_needed
CUTLASS_HOST_DEVICE constexpr bool is_source_needed() const
Returns true if source is needed based on state of runtime arguments.
Definition: conversion_op.h:89
cutlass::epilogue::thread::Convert::kRound
static FloatRoundStyle const kRound
Definition: conversion_op.h:66
cutlass::epilogue::thread::Convert::ElementAccumulator
ElementAccumulator_ ElementAccumulator
Definition: conversion_op.h:57
array.h
Statically sized array of elements that accommodates all CUTLASS-supported numeric types and is safe ...
cutlass::epilogue::thread::Convert::operator()
CUTLASS_HOST_DEVICE FragmentOutput operator()(FragmentAccumulator const &accumulator, FragmentOutput const &source, ElementCompute uniform=ElementCompute(0)) const
Computes linear scaling: D = alpha * accumulator + beta * source.
Definition: conversion_op.h:102
numeric_conversion.h
Boost-like numeric conversion operator for CUTLASS numeric types.
cutlass::epilogue::thread::Convert::kCount
static int const kCount
Definition: conversion_op.h:60
CUTLASS_HOST_DEVICE
#define CUTLASS_HOST_DEVICE
Definition: cutlass.h:89
numeric_types.h
Top-level include for all CUTLASS numeric types.
cutlass::epilogue::thread::Convert::Params::Params
CUTLASS_HOST_DEVICE Params()
Definition: conversion_op.h:76
cutlass::epilogue::thread::Convert::ElementOutput
ElementOutput_ ElementOutput
Definition: conversion_op.h:56
cutlass::epilogue::thread::Convert::FragmentAccumulator
Array< ElementAccumulator, kCount > FragmentAccumulator
Definition: conversion_op.h:63
cutlass::FloatRoundStyle::round_to_nearest
round to nearest even
cutlass::epilogue::thread::Convert::ElementCompute
ElementAccumulator_ ElementCompute
Definition: conversion_op.h:58
cutlass::FloatRoundStyle
FloatRoundStyle
Definition: numeric_conversion.h:43
cutlass::epilogue::thread::Convert::is_source_ever_needed
CUTLASS_HOST_DEVICE constexpr bool is_source_ever_needed() const
Definition: conversion_op.h:96
cutlass::NumericArrayConverter
Conversion operator for Array.
Definition: numeric_conversion.h:294
cutlass::epilogue::thread::Convert::FragmentOutput
Array< ElementOutput, kCount > FragmentOutput
Definition: conversion_op.h:62
cutlass.h
Basic include for CUTLASS.
cutlass::epilogue::thread::Convert::Params
Host-constructable parameters structure.
Definition: conversion_op.h:69
functional.h
Define basic numeric operators with specializations for Array<T, N>. SIMD-ize where possible...