| CUTLASS CUDA Templates for Linear Algebra Subroutines and Solvers

- Main Page
- Modules
- Namespaces
- Classes
- Files

- File List
- File Members

conversion_op.h
Go to the documentation of this file.
1 /***************************************************************************************************
2 * Copyright (c) 2017-2019, NVIDIA CORPORATION. All rights reserved.
3 *
4 * Redistribution and use in source and binary forms, with or without modification, are permitted
5 * provided that the following conditions are met:
6 * * Redistributions of source code must retain the above copyright notice, this list of
7 * conditions and the following disclaimer.
8 * * Redistributions in binary form must reproduce the above copyright notice, this list of
9 * conditions and the following disclaimer in the documentation and/or other materials
10 * provided with the distribution.
11 * * Neither the name of the NVIDIA CORPORATION nor the names of its contributors may be used
12 * to endorse or promote products derived from this software without specific prior written
13 * permission.
14 *
15 * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND ANY EXPRESS OR
16 * IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND
17 * FITNESS FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE
18 * FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING,
19 * BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS;
20 * OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT,
21 * STRICT LIABILITY, OR TOR (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
22 * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
23 *
24 **************************************************************************************************/
29 #pragma once
30
31 #include "cutlass/cutlass.h"
32 #include "cutlass/numeric_types.h"
33 #include "cutlass/array.h"
34 #include "cutlass/functional.h"
35 #include "cutlass/numeric_conversion.h"
36
38
39 namespace cutlass {
40 namespace epilogue {
41 namespace thread {
42
44
47 template <
48 typename ElementOutput_,
49 int Count,
50 typename ElementAccumulator_ = ElementOutput_,
51 FloatRoundStyle Round = FloatRoundStyle::round_to_nearest
52 >
53 class Convert {
54 public:
55
56 using ElementOutput = ElementOutput_;
57 using ElementAccumulator = ElementAccumulator_;
58 using ElementCompute = ElementAccumulator_;
59
60 static int const kCount = Count;
61
62 using FragmentOutput = Array<ElementOutput, kCount>;
63 using FragmentAccumulator = Array<ElementAccumulator, kCount>;
64 using ComputeFragment = FragmentAccumulator;
65
66 static FloatRoundStyle const kRound = Round;
67
69 struct Params {
70
71 //
72 // Methods
73 //
74
75 CUTLASS_HOST_DEVICE
76 Params() {}
77 };
78
79 public:
80
82 CUTLASS_HOST_DEVICE
83 Convert(Params const &params = Params()) {
84
85 }
86
88 CUTLASS_HOST_DEVICE
89 constexpr bool is_source_needed() const {
90 return false;
91 }
92
95 CUTLASS_HOST_DEVICE
96 constexpr bool is_source_ever_needed() const {
97 return false;
98 }
99
101 CUTLASS_HOST_DEVICE
102 FragmentOutput operator()(
103 FragmentAccumulator const &accumulator,
104 FragmentOutput const &source,
105 ElementCompute uniform = ElementCompute(0)) const {
106
107 // Convert to destination numeric type
108 NumericArrayConverter<ElementOutput, ElementAccumulator, kCount, Round> destination_converter;
109
110 return destination_converter(accumulator);
111 }
112 };
113
115
116 } // namespace thread
117 } // namespace epilogue
118 } // 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...