Back to Cutlass

CUTLASS: conversion_op.h Source File

docs/conversion__op_8h_source.html

4.4.213.3 KB
Original Source

| | CUTLASS

CUDA Templates for Linear Algebra Subroutines and Solvers |

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 <

48typename ElementOutput_,

49int Count,

50typename ElementAccumulator_ = ElementOutput_,

51FloatRoundStyle Round = FloatRoundStyle::round_to_nearest

52 >

53 class Convert {

54 public:

55

56using ElementOutput = ElementOutput_;

57using ElementAccumulator = ElementAccumulator_;

58using ElementCompute = ElementAccumulator_;

59

60static int const kCount = Count;

61

62using FragmentOutput = Array<ElementOutput, kCount>;

63using FragmentAccumulator = Array<ElementAccumulator, kCount>;

64using ComputeFragment = FragmentAccumulator;

65

66static FloatRoundStyle const kRound = Round;

67

69struct Params {

70

71//

72// Methods

73//

74

75CUTLASS_HOST_DEVICE

76Params() {}

77 };

78

79 public:

80

82CUTLASS_HOST_DEVICE

83Convert(Params const &params = Params()) {

84

85 }

86

88CUTLASS_HOST_DEVICE

89constexpr bool is_source_needed() const {

90return false;

91 }

92

95CUTLASS_HOST_DEVICE

96constexpr bool is_source_ever_needed() const {

97return false;

98 }

99

101CUTLASS_HOST_DEVICE

102FragmentOutput operator()(

103FragmentAccumulator const &accumulator,

104FragmentOutput const &source,

105ElementCompute uniform = ElementCompute(0)) const {

106

107// Convert to destination numeric type

108NumericArrayConverter<ElementOutput, ElementAccumulator, kCount, Round> destination_converter;

109

110return 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...


Generated by 1.8.11