Cutlass CUDA Templates for Linear Algebra Subroutines and Solvers

coord.h
Go to the documentation of this file.
1 /***************************************************************************************************
2 * Copyright (c) 2017-2018, 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/util/platform.h"
33
34 namespace cutlass {
35
37
39 struct Identity {
42 enum Kind { Additive = 0, Multiplicative = 1 };
43 };
44
46
48 template <int Rank_, typename Index_ = int>
49 struct Coord {
50 //
51 // Type and constant definitions
52 //
53
55 static int const kRank = Rank_;
56
58 static int const N = Rank_;
59
61 typedef Index_ Index;
62
63 //
64 // Data members
65 //
66
68 Index idx[kRank];
69
70 //
71 // Methods
72 //
73
75 CUTLASS_HOST_DEVICE
76 Coord(Index value = 0) {
77 for (int i = 0; i < kRank; ++i) {
78 idx[i] = value;
79 }
80 }
81
83 CUTLASS_HOST_DEVICE
84 Coord(Index _idx[]) {
85 for (int i = 0; i < kRank; ++i) {
86 idx[i] = _idx[i];
87 }
88 }
89
91 CUTLASS_HOST_DEVICE
92 Coord(Coord<kRank> const &coord) {
93 for (int i = 0; i < kRank; ++i) {
94 idx[i] = coord[i];
95 }
96 }
97
100 template <int Slice>
101 CUTLASS_HOST_DEVICE
102 Coord<Slice> slice(int start = 0, Index identity = 0) const {
103 Coord<Slice> result;
104 for (int i = 0; i < Slice; ++i) {
105 if (i + start < kRank) {
106 slice[i] = idx[i + start];
107 }
108 else {
109 slice[i] = identity;
110 }
111 }
112 return result;
113 }
114
116 CUTLASS_HOST_DEVICE
117 operator bool() const {
118 for (int i = 0; i < kRank; ++i) {
119 if (idx[i]) {
120 return true;
121 }
122 }
123 return false;
124 }
125
127 CUTLASS_HOST_DEVICE
128 bool operator!() const {
129 for (int i = 0; i < kRank; ++i) {
130 if (idx[i]) {
131 return false;
132 }
133 }
134 return true;
135 }
136
138 CUTLASS_HOST_DEVICE
139 Coord operator+(Coord const& b) const {
140 Coord c;
141 for (int i = 0; i < kRank; ++i) {
142 c.idx[i] = idx[i] + b.idx[i];
143 }
144 return c;
145 }
146
148 CUTLASS_HOST_DEVICE
149 Coord operator-(Coord const& b) const {
150 Coord c;
151 for (int i = 0; i < kRank; ++i) {
152 c.idx[i] = idx[i] - b.idx[i];
153 }
154 return c;
155 }
156
158 CUTLASS_HOST_DEVICE
159 Coord operator*(Coord const& b) const {
160 Coord c;
161 for (int i = 0; i < kRank; ++i) {
162 c.idx[i] = idx[i] * b.idx[i];
163 }
164 return c;
165 }
166
168 CUTLASS_HOST_DEVICE
169 Coord operator/(Coord const& b) const {
170 Coord c;
171 for (int i = 0; i < kRank; ++i) {
172 c.idx[i] = idx[i] / b.idx[i];
173 }
174 return c;
175 }
176
178 CUTLASS_HOST_DEVICE
179 Coord& operator+=(Coord const& b) {
180 for (int i = 0; i < kRank; ++i) {
181 idx[i] += b.idx[i];
182 }
183 return *this;
184 }
185
187 CUTLASS_HOST_DEVICE
188 Coord& operator-=(Coord const& b) {
189 for (int i = 0; i < kRank; ++i) {
190 idx[i] -= b.idx[i];
191 }
192 return *this;
193 }
194
196 CUTLASS_HOST_DEVICE
197 Coord& operator*=(Coord const& b) {
198 for (int i = 0; i < kRank; ++i) {
199 idx[i] *= b.idx[i];
200 }
201 return *this;
202 }
203
205 CUTLASS_HOST_DEVICE
206 Coord& operator/=(Coord const& b) {
207 for (int i = 0; i < kRank; ++i) {
208 idx[i] /= b.idx[i];
209 }
210 return *this;
211 }
212
214 CUTLASS_HOST_DEVICE Index& operator[](int dim) { return idx[dim]; }
215
217 CUTLASS_HOST_DEVICE Index const& operator[](int dim) const { return idx[dim]; }
218
220 template <typename T>
221 CUTLASS_HOST_DEVICE T dot(Coord const& b, T sum) const {
222 for (int i = 0; i < kRank; ++i) {
223 sum += idx[i] * b.idx[i];
224 }
225 return sum;
226 }
227
229 template <typename T>
230 CUTLASS_HOST_DEVICE T dot(Coord const& b) const {
231 T sum = T(0);
232 for (int i = 0; i < kRank; ++i) {
233 sum += idx[i] * b.idx[i];
234 }
235 return sum;
236 }
237
239 template <int Dim>
240 CUTLASS_HOST_DEVICE Index& at() {
241 return idx[Dim];
242 }
243
245 CUTLASS_HOST_DEVICE
246 Index& at(int dim) { return idx[dim]; }
247
249 template <int Dim>
250 CUTLASS_HOST_DEVICE Index const& at() const {
251 return idx[Dim];
252 }
253
255 CUTLASS_HOST_DEVICE
256 Index const& at(int dim) const { return idx[dim]; }
257
259 CUTLASS_HOST_DEVICE
260 bool operator==(Coord<kRank> const& b) const {
261 bool equal = true;
262 for (int i = 0; equal && i < kRank; ++i) {
263 equal = (idx[i] == b.idx[i]);
264 }
265 return equal;
266 }
267
269 CUTLASS_HOST_DEVICE
270 bool operator!=(Coord<kRank> const& b) const { return !(*this == b); }
271
273 CUTLASS_HOST_DEVICE
274 Coord& clamp(Coord<kRank> const& max, Coord<kRank> const& min = Coord<kRank>()) {
275 for (int i = 0; i < kRank; ++i) {
276 idx[i] = __NV_STD_MAX(__NV_STD_MIN(idx[i], max.idx[i]), min.idx[i]);
277 }
278 return *this;
279 }
280
282 CUTLASS_HOST_DEVICE
283 Index count() const {
284 Index product = idx[0];
285 for (int i = 1; i < kRank; ++i) {
286 product *= idx[i];
287 }
288 return product;
289 }
290
292 CUTLASS_HOST_DEVICE
293 bool operator<(Coord<kRank> const &b) const {
294 for (int i = 0; i < kRank; ++i) {
295 if (!(idx[i] < b[i])) {
296 return false;
297 }
298 }
299 return true;
300 }
301
303 CUTLASS_HOST_DEVICE
304 bool operator<=(Coord<kRank> const &b) const {
305 for (int i = 0; i < kRank; ++i) {
306 if (!(idx[i] <= b[i])) {
307 return false;
308 }
309 }
310 return true;
311 }
312 };
313
315
317 template <typename T, int Rank, typename Index>
318 CUTLASS_HOST_DEVICE
319 Coord<Rank, Index> operator*(T s, Coord<Rank, Index> coord) {
320 CUTLASS_PRAGMA_UNROLL
321 for (int i = 0; i < Rank; ++i) {
322 coord[i] *= s;
323 }
324 return coord;
325 }
326
328 template <typename T, int Rank, typename Index>
329 CUTLASS_HOST_DEVICE
330 Coord<Rank, Index> operator*(Coord<Rank, Index> coord, T s) {
331 CUTLASS_PRAGMA_UNROLL
332 for (int i = 0; i < Rank; ++i) {
333 coord[i] *= s;
334 }
335 return coord;
336 }
337
339 template <typename T, int Rank, typename Index>
340 CUTLASS_HOST_DEVICE
341 Coord<Rank, Index> operator/(T s, Coord<Rank, Index> coord) {
342 CUTLASS_PRAGMA_UNROLL
343 for (int i = 0; i < Rank; ++i) {
344 coord[i] = s / coord[i];
345 }
346 return coord;
347 }
348
350 template <typename T, int Rank, typename Index>
351 CUTLASS_HOST_DEVICE
352 Coord<Rank, Index> operator/(Coord<Rank, Index> coord, T s) {
353 CUTLASS_PRAGMA_UNROLL
354 for (int i = 0; i < Rank; ++i) {
355 coord[i] /= s;
356 }
357 return coord;
358 }
359
361 //
362 // Integer-valued make_Coord
363 //
365
367 CUTLASS_HOST_DEVICE
368 Coord<1> make_Coord(int _0) {
369 int values[1] = {_0};
370 return Coord<1>(values);
371 }
372
374 CUTLASS_HOST_DEVICE
375 Coord<2> make_Coord(int _0, int _1) {
376 int values[2] = {_0, _1};
377 return Coord<2>(values);
378 }
379
381 CUTLASS_HOST_DEVICE
382 Coord<3> make_Coord(int _0, int _1, int _2) {
383 int values[3] = {_0, _1, _2};
384 return Coord<3>(values);
385 }
386
388 CUTLASS_HOST_DEVICE
389 Coord<4> make_Coord(int _0, int _1, int _2, int _3) {
390 int values[4] = {_0, _1, _2, _3};
391 return Coord<4>(values);
392 }
393
395
396 template <typename Shape_>
397 CUTLASS_HOST_DEVICE Coord<3> make_Coord_from_shape() {
398 return make_Coord(Shape_::kD, Shape_::kH, Shape_::kW);
399 }
400
402
403 } // namespace cutlass
cutlass::Identity
Describes identity elements.
Definition: coord.h:39
cutlass::platform::max
CUTLASS_HOST_DEVICE constexpr const T & max(const T &a, const T &b)
std::max
Definition: platform.h:215
CUTLASS_PRAGMA_UNROLL
#define CUTLASS_PRAGMA_UNROLL
Definition: performance_tuning.h:35
cutlass
Definition: convert.h:33
cutlass::Coord::operator-
CUTLASS_HOST_DEVICE Coord operator-(Coord const &b) const
Element-wise subtraction.
Definition: coord.h:149
cutlass::operator*
CUTLASS_HOST_DEVICE Coord< Rank, Index > operator*(T s, Coord< Rank, Index > coord)
Scalar multiplication.
Definition: coord.h:319
cutlass::Coord::at
CUTLASS_HOST_DEVICE Index const & at(int dim) const
Access via index; may limit unrolling potential.
Definition: coord.h:256
cutlass::Coord::operator[]
CUTLASS_HOST_DEVICE Index const & operator[](int dim) const
Member access operator.
Definition: coord.h:217
cutlass::Coord::operator/
CUTLASS_HOST_DEVICE Coord operator/(Coord const &b) const
Element-wise division.
Definition: coord.h:169
cutlass::Coord::operator[]
CUTLASS_HOST_DEVICE Index & operator[](int dim)
Member access operator.
Definition: coord.h:214
cutlass::make_Coord
CUTLASS_HOST_DEVICE Coord< 1 > make_Coord(int _0)
Helper to make a 2-element coordinate.
Definition: coord.h:368
cutlass::Coord::kRank
static int const kRank
Number of elements in Coord.
Definition: coord.h:55
cutlass::Coord::Index
Index_ Index
Index type used to store elements.
Definition: coord.h:61
cutlass::Coord::operator*=
CUTLASS_HOST_DEVICE Coord & operator*=(Coord const &b)
In-place multiplication.
Definition: coord.h:197
cutlass::Coord::at
CUTLASS_HOST_DEVICE Index & at(int dim)
Access via index; may limit unrolling potential.
Definition: coord.h:246
platform.h
C++ features that may be otherwise unimplemented for CUDA device functions.
cutlass::Coord::count
CUTLASS_HOST_DEVICE Index count() const
Returns the product of all elements.
Definition: coord.h:283
cutlass::Coord::operator*
CUTLASS_HOST_DEVICE Coord operator*(Coord const &b) const
Element-wise multiplication.
Definition: coord.h:159
cutlass::Identity::Kind
Kind
Definition: coord.h:42
cutlass::make_Coord_from_shape
CUTLASS_HOST_DEVICE Coord< 3 > make_Coord_from_shape()
Definition: coord.h:397
cutlass::Coord::operator==
CUTLASS_HOST_DEVICE bool operator==(Coord< kRank > const &b) const
Determines if two Coord<> objects are equal.
Definition: coord.h:260
cutlass::Coord::N
static int const N
Number of elements in Coord, aliased for compatibility.
Definition: coord.h:58
__NV_STD_MAX
#define __NV_STD_MAX(a, b)
Select maximum(a, b)
Definition: platform.h:163
cutlass::Coord::idx
Index idx[kRank]
Indices.
Definition: coord.h:68
__NV_STD_MIN
#define __NV_STD_MIN(a, b)
Select minimum(a, b)
Definition: platform.h:168
cutlass::Coord::operator-=
CUTLASS_HOST_DEVICE Coord & operator-=(Coord const &b)
In-place subtraction.
Definition: coord.h:188
cutlass::Coord::operator+=
CUTLASS_HOST_DEVICE Coord & operator+=(Coord const &b)
In-place addition.
Definition: coord.h:179
CUTLASS_HOST_DEVICE
#define CUTLASS_HOST_DEVICE
Definition: cutlass.h:46
cutlass::Coord::operator!=
CUTLASS_HOST_DEVICE bool operator!=(Coord< kRank > const &b) const
Not equal.
Definition: coord.h:270
cutlass::platform::min
CUTLASS_HOST_DEVICE constexpr const T & min(const T &a, const T &b)
std::min
Definition: platform.h:209
cutlass::Coord::at
CUTLASS_HOST_DEVICE Index & at()
Gets the index of a given Coord element.
Definition: coord.h:240
cutlass::Coord::operator/=
CUTLASS_HOST_DEVICE Coord & operator/=(Coord const &b)
In-place division.
Definition: coord.h:206
cutlass::Identity::Multiplicative
Definition: coord.h:42
cutlass::Coord::slice
CUTLASS_HOST_DEVICE Coord< Slice > slice(int start=0, Index identity=0) const
Definition: coord.h:102
cutlass::Coord
Statically-sized array specifying Coords within a tensor.
Definition: coord.h:49
cutlass::operator/
CUTLASS_HOST_DEVICE Coord< Rank, Index > operator/(T s, Coord< Rank, Index > coord)
Scalar division.
Definition: coord.h:341
cutlass::Coord::at
CUTLASS_HOST_DEVICE Index const & at() const
Gets the index of a given Coord element.
Definition: coord.h:250
cutlass::Coord::dot
CUTLASS_HOST_DEVICE T dot(Coord const &b, T sum) const
Computes the dot product of two Coord instances.
Definition: coord.h:221
cutlass::Coord::Coord
CUTLASS_HOST_DEVICE Coord(Index value=0)
Default ctor initializes uniformly.
Definition: coord.h:76
cutlass::Identity::Additive
Definition: coord.h:42
cutlass::Coord::clamp
CUTLASS_HOST_DEVICE Coord & clamp(Coord< kRank > const &max, Coord< kRank > const &min=Coord< kRank >())
Clamps a coordinate to a range specified by maximum and minimum values.
Definition: coord.h:274
cutlass::Coord::Coord
CUTLASS_HOST_DEVICE Coord(Index _idx[])
Constructs from an array of integers.
Definition: coord.h:84
cutlass::Coord::dot
CUTLASS_HOST_DEVICE T dot(Coord const &b) const
Computes the dot product of two Coord instances.
Definition: coord.h:230
cutlass::Coord::operator+
CUTLASS_HOST_DEVICE Coord operator+(Coord const &b) const
Element-wise addition.
Definition: coord.h:139
cutlass.h
Basic include for CUTLASS macros.
cutlass::Coord::Coord
CUTLASS_HOST_DEVICE Coord(Coord< kRank > const &coord)
Constructs from an array of integers.
Definition: coord.h:92
cutlass::Coord::operator!
CUTLASS_HOST_DEVICE bool operator!() const
Returns true if Coord is uniformly zero.
Definition: coord.h:128