docs/pitch__linear_8h_source.html
| | CUTLASS
CUDA Templates for Linear Algebra Subroutines and Solvers |
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 <
40int Contiguous,
41int Strided
42 >
43 struct PitchLinearShape {
44static int const kContiguous = Contiguous;
45static int const kStrided = Strided;
46static int const kCount = Contiguous * Strided;
47 };
48
50
52 struct PitchLinearCoord : public Coord<2, int> {
53 public:
54
57
59using Base = Coord<2, Index>;
60
61 private:
62
64static int const kContiguous = 0;
65
67static int const kStrided = 1;
68
69 public:
70
71//
72// Methods
73//
74
77PitchLinearCoord() { }
78
81PitchLinearCoord(Coord<2, Index> const &coord): Base(coord) { }
82
85PitchLinearCoord(Index contiguous_, Index strided_): Base(make_Coord(contiguous_, strided_)) { }
86
89Index const & contiguous() const { return this->at(kContiguous); }
90
93Index & contiguous() { return this->at(kContiguous); }
94
97Index const & strided() const { return this->at(kStrided); }
98
101Index & strided() { return this->at(kStrided); }
102
103//
104// Coord operators
105//
106
109PitchLinearCoord operator+(Base const& b) const {
110return PitchLinearCoord(Base::operator+(b));
111 }
112
115PitchLinearCoord operator-(Base const& b) const {
116return PitchLinearCoord(Base::operator-(b));
117 }
118
121PitchLinearCoord operator*(Base const& b) const {
122return PitchLinearCoord(Base::operator*(b));
123 }
124
127PitchLinearCoord operator/(Base const& b) const {
128return PitchLinearCoord(Base::operator/(b));
129 }
130
133PitchLinearCoord& operator+=(Base const& b) {
134Base::operator+=(b);
135return *this;
136 }
137
140PitchLinearCoord& operator-=(Base const& b) {
141Base::operator-=(b);
142return *this;
143 }
144
147PitchLinearCoord& operator*=(Base const& b) {
148Base::operator*=(b);
149return *this;
150 }
151
154PitchLinearCoord& operator/=(Base const& b) {
155Base::operator/=(b);
156return *this;
157 }
158 };
159
161
163 class PitchLinear {
164 public:
166static int const kRank = 2;
167
169static int const kStrideRank = 1;
170
173
176
178using TensorCoord = PitchLinearCoord;
179
181using Stride = Coord<kStrideRank, Index>;
182
183 private:
184//
185// Data members
186//
187
189Stride stride_;
190
191 public:
192//
193// Methods
194//
195
198PitchLinear(Index ldm = 0): stride_(ldm) { }
199
202PitchLinear(Stride _stride): stride_(_stride) { }
203
206static PitchLinear packed(TensorCoord const &extent) {
207return PitchLinear(extent.contiguous());
208 }
209
213LongIndex operator()(TensorCoord const &coord) const {
214return coord.contiguous() + coord.strided() * stride_[0];
215 }
216
219TensorCoord inverse(LongIndex index) const {
220return make_Coord(
221Index(index % stride_[0]),
222Index(index / stride_[0])
223 );
224 }
225
229return stride_;
230 }
231
235return stride_;
236 }
237
240Index stride(int rank) const {
241return stride_[rank];
242 }
243
247return stride_[rank];
248 }
249
252LongIndex capacity(TensorCoord const &extent) const {
253return 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
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
Mapping function for pitch-linear memory.
Definition: pitch_linear.h:163
A Coord is a coordinate of arbitrary rank into a tensor or matrix.
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_HOST_DEVICE half_t & operator/=(half_t &lhs, half_t const &rhs)
Definition: half.h:684
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_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
#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_HOST_DEVICE half_t & operator*=(half_t &lhs, half_t const &rhs)
Definition: half.h:674
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
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
Generated by 1.8.11