| CUTLASS CUDA Templates for Linear Algebra Subroutines and Solvers

- Main Page
- Modules
- Namespaces
- Classes
- Files

- File List
- File Members

pitch_linear.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 **************************************************************************************************/
28 #pragma once
29
30 #include "cutlass/cutlass.h"
31 #include "cutlass/coord.h"
32
33 namespace cutlass {
34 namespace layout {
35
37
39 template <
40 int Contiguous,
41 int Strided
42 >
43 struct PitchLinearShape {
44 static int const kContiguous = Contiguous;
45 static int const kStrided = Strided;
46 static int const kCount = Contiguous * Strided;
47 };
48
50
52 struct PitchLinearCoord : public Coord<2, int> {
53 public:
54
56 using Index = int;
57
59 using Base = Coord<2, Index>;
60
61 private:
62
64 static int const kContiguous = 0;
65
67 static int const kStrided = 1;
68
69 public:
70
71 //
72 // Methods
73 //
74
76 CUTLASS_HOST_DEVICE
77 PitchLinearCoord() { }
78
80 CUTLASS_HOST_DEVICE
81 PitchLinearCoord(Coord<2, Index> const &coord): Base(coord) { }
82
84 CUTLASS_HOST_DEVICE
85 PitchLinearCoord(Index contiguous_, Index strided_): Base(make_Coord(contiguous_, strided_)) { }
86
88 CUTLASS_HOST_DEVICE
89 Index const & contiguous() const { return this->at(kContiguous); }
90
92 CUTLASS_HOST_DEVICE
93 Index & contiguous() { return this->at(kContiguous); }
94
96 CUTLASS_HOST_DEVICE
97 Index const & strided() const { return this->at(kStrided); }
98
100 CUTLASS_HOST_DEVICE
101 Index & strided() { return this->at(kStrided); }
102
103 //
104 // Coord operators
105 //
106
108 CUTLASS_HOST_DEVICE
109 PitchLinearCoord operator+(Base const& b) const {
110 return PitchLinearCoord(Base::operator+(b));
111 }
112
114 CUTLASS_HOST_DEVICE
115 PitchLinearCoord operator-(Base const& b) const {
116 return PitchLinearCoord(Base::operator-(b));
117 }
118
120 CUTLASS_HOST_DEVICE
121 PitchLinearCoord operator*(Base const& b) const {
122 return PitchLinearCoord(Base::operator*(b));
123 }
124
126 CUTLASS_HOST_DEVICE
127 PitchLinearCoord operator/(Base const& b) const {
128 return PitchLinearCoord(Base::operator/(b));
129 }
130
132 CUTLASS_HOST_DEVICE
133 PitchLinearCoord& operator+=(Base const& b) {
134 Base::operator+=(b);
135 return *this;
136 }
137
139 CUTLASS_HOST_DEVICE
140 PitchLinearCoord& operator-=(Base const& b) {
141 Base::operator-=(b);
142 return *this;
143 }
144
146 CUTLASS_HOST_DEVICE
147 PitchLinearCoord& operator*=(Base const& b) {
148 Base::operator*=(b);
149 return *this;
150 }
151
153 CUTLASS_HOST_DEVICE
154 PitchLinearCoord& operator/=(Base const& b) {
155 Base::operator/=(b);
156 return *this;
157 }
158 };
159
161
163 class PitchLinear {
164 public:
166 static int const kRank = 2;
167
169 static int const kStrideRank = 1;
170
172 using Index = int32_t;
173
175 using LongIndex = int64_t;
176
178 using TensorCoord = PitchLinearCoord;
179
181 using Stride = Coord<kStrideRank, Index>;
182
183 private:
184 //
185 // Data members
186 //
187
189 Stride stride_;
190
191 public:
192 //
193 // Methods
194 //
195
197 CUTLASS_HOST_DEVICE
198 PitchLinear(Index ldm = 0): stride_(ldm) { }
199
201 CUTLASS_HOST_DEVICE
202 PitchLinear(Stride _stride): stride_(_stride) { }
203
205 CUTLASS_HOST_DEVICE
206 static PitchLinear packed(TensorCoord const &extent) {
207 return PitchLinear(extent.contiguous());
208 }
209
212 CUTLASS_HOST_DEVICE
213 LongIndex operator()(TensorCoord const &coord) const {
214 return coord.contiguous() + coord.strided() * stride_[0];
215 }
216
218 CUTLASS_HOST_DEVICE
219 TensorCoord inverse(LongIndex index) const {
220 return make_Coord(
221 Index(index % stride_[0]),
222 Index(index / stride_[0])
223 );
224 }
225
227 CUTLASS_HOST_DEVICE
228 Stride stride() const {
229 return stride_;
230 }
231
233 CUTLASS_HOST_DEVICE
234 Stride & stride() {
235 return stride_;
236 }
237
239 CUTLASS_HOST_DEVICE
240 Index stride(int rank) const {
241 return stride_[rank];
242 }
243
245 CUTLASS_HOST_DEVICE
246 Index & stride(int rank) {
247 return stride_[rank];
248 }
249
251 CUTLASS_HOST_DEVICE
252 LongIndex capacity(TensorCoord const &extent) const {
253 return extent.strided() * stride_[0];
254 }
255 };
256
258
259 } // namespace layout
260 } // namespace cutlass
261
cutlass::layout::PitchLinearCoord::Index
int Index
Integer-valued index.
Definition: pitch_linear.h:56
cutlass::layout::PitchLinearShape::kCount
static int const kCount
Definition: pitch_linear.h:46
cutlass
Definition: aligned_buffer.h:35
cutlass::layout::PitchLinearCoord
Coordinate in pitch-linear space.
Definition: pitch_linear.h:52
cutlass::layout::PitchLinearCoord::operator-
CUTLASS_HOST_DEVICE PitchLinearCoord operator-(Base const &b) const
Element-wise subtraction.
Definition: pitch_linear.h:115
cutlass::layout::PitchLinear
Mapping function for pitch-linear memory.
Definition: pitch_linear.h:163
coord.h
A Coord is a coordinate of arbitrary rank into a tensor or matrix.
cutlass::make_Coord
CUTLASS_HOST_DEVICE Coord< 1 > make_Coord(int _0)
Helper to make a 2-element coordinate.
Definition: coord.h:387
cutlass::layout::PitchLinear::stride
CUTLASS_HOST_DEVICE Stride stride() const
Returns the stride of the layout.
Definition: pitch_linear.h:228
cutlass::layout::PitchLinear::packed
static CUTLASS_HOST_DEVICE PitchLinear packed(TensorCoord const &extent)
Helper returns a layout to a tightly packed tensor.
Definition: pitch_linear.h:206
cutlass::operator/=
CUTLASS_HOST_DEVICE half_t & operator/=(half_t &lhs, half_t const &rhs)
Definition: half.h:684
cutlass::operator+=
CUTLASS_HOST_DEVICE half_t & operator+=(half_t &lhs, half_t const &rhs)
Definition: half.h:654
cutlass::layout::PitchLinearShape
Template defining a shape used by pitch-linear operators.
Definition: pitch_linear.h:43
cutlass::operator-=
CUTLASS_HOST_DEVICE half_t & operator-=(half_t &lhs, half_t const &rhs)
Definition: half.h:664
cutlass::layout::PitchLinear::LongIndex
int64_t LongIndex
Long index type used for offsets.
Definition: pitch_linear.h:175
cutlass::layout::PitchLinearCoord::operator+=
CUTLASS_HOST_DEVICE PitchLinearCoord & operator+=(Base const &b)
In-place addition.
Definition: pitch_linear.h:133
cutlass::layout::PitchLinear::PitchLinear
CUTLASS_HOST_DEVICE PitchLinear(Stride _stride)
Constructor.
Definition: pitch_linear.h:202
cutlass::layout::PitchLinearShape::kStrided
static int const kStrided
Definition: pitch_linear.h:45
cutlass::layout::PitchLinearCoord::operator*=
CUTLASS_HOST_DEVICE PitchLinearCoord & operator*=(Base const &b)
In-place multiplication.
Definition: pitch_linear.h:147
cutlass::layout::PitchLinearShape::kContiguous
static int const kContiguous
Definition: pitch_linear.h:44
cutlass::layout::PitchLinear::inverse
CUTLASS_HOST_DEVICE TensorCoord inverse(LongIndex index) const
Returns the logical coordinate given an offset.
Definition: pitch_linear.h:219
cutlass::layout::PitchLinear::stride
CUTLASS_HOST_DEVICE Index & stride(int rank)
Returns the stride of the layout.
Definition: pitch_linear.h:246
cutlass::layout::PitchLinear::stride
CUTLASS_HOST_DEVICE Stride & stride()
Returns the stride of the layout.
Definition: pitch_linear.h:234
CUTLASS_HOST_DEVICE
#define CUTLASS_HOST_DEVICE
Definition: cutlass.h:89
cutlass::layout::PitchLinearCoord::contiguous
CUTLASS_HOST_DEVICE Index const & contiguous() const
Returns the contiguous dimension.
Definition: pitch_linear.h:89
cutlass::layout::PitchLinearCoord::PitchLinearCoord
CUTLASS_HOST_DEVICE PitchLinearCoord(Index contiguous_, Index strided_)
Helper to construct from a row and column.
Definition: pitch_linear.h:85
cutlass::layout::PitchLinear::Index
int32_t Index
Index type used for coordinates.
Definition: pitch_linear.h:172
cutlass::operator*=
CUTLASS_HOST_DEVICE half_t & operator*=(half_t &lhs, half_t const &rhs)
Definition: half.h:674
cutlass::Coord
Statically-sized array specifying Coords within a tensor.
Definition: coord.h:43
cutlass::layout::PitchLinear::operator()
CUTLASS_HOST_DEVICE LongIndex operator()(TensorCoord const &coord) const
Definition: pitch_linear.h:213
cutlass::layout::PitchLinearCoord::operator/=
CUTLASS_HOST_DEVICE PitchLinearCoord & operator/=(Base const &b)
In-place division.
Definition: pitch_linear.h:154
cutlass::layout::PitchLinear::stride
CUTLASS_HOST_DEVICE Index stride(int rank) const
Returns the stride of the layout.
Definition: pitch_linear.h:240
cutlass::layout::PitchLinear::PitchLinear
CUTLASS_HOST_DEVICE PitchLinear(Index ldm=0)
Constructor.
Definition: pitch_linear.h:198
cutlass::layout::PitchLinearCoord::strided
CUTLASS_HOST_DEVICE Index & strided()
Returns the column of the coordinate.
Definition: pitch_linear.h:101
cutlass::layout::PitchLinearCoord::operator/
CUTLASS_HOST_DEVICE PitchLinearCoord operator/(Base const &b) const
Element-wise division.
Definition: pitch_linear.h:127
cutlass::layout::PitchLinearCoord::operator*
CUTLASS_HOST_DEVICE PitchLinearCoord operator*(Base const &b) const
Element-wise multiplication.
Definition: pitch_linear.h:121
cutlass::layout::PitchLinearCoord::PitchLinearCoord
CUTLASS_HOST_DEVICE PitchLinearCoord(Coord< 2, Index > const &coord)
Constructs from Coord<2>
Definition: pitch_linear.h:81
cutlass::layout::PitchLinearCoord::operator+
CUTLASS_HOST_DEVICE PitchLinearCoord operator+(Base const &b) const
Element-wise addition.
Definition: pitch_linear.h:109
cutlass::layout::PitchLinearCoord::contiguous
CUTLASS_HOST_DEVICE Index & contiguous()
Returns the contiguous dimension.
Definition: pitch_linear.h:93
cutlass::layout::PitchLinearCoord::PitchLinearCoord
CUTLASS_HOST_DEVICE PitchLinearCoord()
Default ctor.
Definition: pitch_linear.h:77
cutlass.h
Basic include for CUTLASS.
cutlass::layout::PitchLinearCoord::operator-=
CUTLASS_HOST_DEVICE PitchLinearCoord & operator-=(Base const &b)
In-place subtraction.
Definition: pitch_linear.h:140
cutlass::layout::PitchLinearCoord::strided
CUTLASS_HOST_DEVICE Index const & strided() const
Returns the column of the coordinate.
Definition: pitch_linear.h:97
cutlass::layout::PitchLinear::capacity
CUTLASS_HOST_DEVICE LongIndex capacity(TensorCoord const &extent) const
Compute the number of contiguous elements needed to store a tensor with the given size...
Definition: pitch_linear.h:252