Cutlass CUDA Templates for Linear Algebra Subroutines and Solvers

gemm_global_tile.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 **************************************************************************************************/
28 #pragma once
29
30 #include "cutlass/coord.h"
31 #include "cutlass/util/platform.h"
32
33 #include "cutlass/gemm/gemm_operand.h"
34 #include "cutlass/matrix_traits.h"
35 #include "cutlass/predicate_vector.h"
36 #include "cutlass/reshape_tile.h"
37 #include "cutlass/tile_iterator.h"
38
39 namespace cutlass {
40 namespace gemm {
41
43
44 // The following functor reshapes a tile of threads to match a tile of data. The idea is that when
45 // the user wants to build the iterator traits, he/she may want to specify the tile independently
46 // from the number of scalars loaded/stored per instruction. For example, in the row-major version
47 // with a tile of size 128x8 - the user may want to that the iterator works with 32x8 threads if
48 // each thread loads 1 scalar per LDG. If the user changes to 4 scalars per LDG, then the tile of
49 // threads has to change. The code below detects that and correct the code automatically - it is
50 // a helper when the user does not specify the right configuration.
51
52 template <typename Tile_, typename Threads_, bool = (Tile_::kW < Threads_::kW)>
53 struct ReshapeThreads {
54 typedef Threads_ Threads;
55 };
56
57 template <typename Tile_, typename Threads_>
58 struct ReshapeThreads<Tile_, Threads_, true> {
59 typedef Shape<Threads_::kD, Threads_::kH * Threads_::kW / Tile_::kW, Tile_::kW, 1> Threads;
60 };
61
63
64 template <GemmOperand::Kind kOperand_,
65 MatrixLayout::Kind kLayout_,
66 typename Scalar_,
67 typename Tile_,
68 typename Threads_,
69 int kAccessSize_>
70 struct GemmGlobalTileTraits {
72 static GemmOperand::Kind const kOperand = kOperand_;
74 static MatrixLayout::Kind const kLayout = kLayout_;
76 typedef typename platform::remove_const<Scalar_>::type Scalar;
78 typedef Scalar_* Pointer;
80 static int const kAccessSize = kAccessSize_;
82 static MemorySpace::Kind const kMemorySpace = MemorySpace::kGlobal;
84 typedef Tile_ Tile;
86 typedef typename ReshapeTile<Tile_, kAccessSize_>::Tile VectorizedTile;
88 typedef typename ReshapeThreads<VectorizedTile, Threads_>::Threads Threads;
90 typedef Shape<1, 1, VectorizedTile::kC> ThreadsDelta;
92 typedef Shape<0, Threads::kH, Threads::kW * kAccessSize> Delta;
93
95 typedef Shape<0, 0, Threads::kW * ThreadsDelta::kW, kAccessSize> ImmediateOffsetStrides;
97 typedef Shape<1,
98 VectorizedTile::kH / Threads::kH,
99 VectorizedTile::kW / Threads::kW,
100 VectorizedTile::kC / kAccessSize>
101 Iterations;
102
103 typedef GemmMultiplicandTraits<Tile, kOperand, kLayout> MultiplicandTraits;
104
106 struct ThreadOffset {
107 CUTLASS_HOST_DEVICE
108 Coord<4> operator()() const {
109 int thread_offset_h = threadIdx.x / Threads::kW * ThreadsDelta::kH;
110 int thread_offset_w = threadIdx.x % Threads::kW * ThreadsDelta::kW;
111
112 return make_Coord(0, thread_offset_h, thread_offset_w, 0);
113 }
114 };
115 };
116
118
119 template <typename Scalar_, typename Tile_, typename Threads_, int kStrideH_, int kAccessSize_>
120 struct GemmGlobalTileCdTraits : public GemmGlobalTileTraits<GemmOperand::kC,
121 MatrixLayout::kColumnMajor,
122 Scalar_,
123 Tile_,
124 Threads_,
125 kAccessSize_> {
127 typedef GemmGlobalTileTraits<GemmOperand::kC,
128 MatrixLayout::kColumnMajor,
129 Scalar_,
130 Tile_,
131 Threads_,
132 kAccessSize_>
133 Base;
134
136 static int const kStrideH = kStrideH_;
138 typedef Shape<0, 0, Base::Delta::kW, Base::Delta::kC> Delta;
139
140 typedef typename Base::Iterations Iterations;
141
142 typedef typename Base::Threads Threads;
143
144 typedef typename Base::ThreadsDelta ThreadsDelta;
145
146 typedef typename Base::ImmediateOffsetStrides ImmediateOffsetStrides;
147
149 struct ThreadOffset {
150 CUTLASS_HOST_DEVICE
151 Coord<4> operator()() const {
152 int thread_offset_h = threadIdx.x / Threads::kW * kStrideH * Iterations::kH;
153 int thread_offset_w = threadIdx.x % Threads::kW * ThreadsDelta::kW;
154
155 return make_Coord(0, thread_offset_h, thread_offset_w, 0);
156 }
157 };
158 };
159
161
162 template <typename TileTraits_, typename Index_ = int>
163 struct GemmGlobalIteratorAb
164 : public TileLoadIterator<TileTraits_,
165 typename TileTraits_::Scalar,
166 TileTraits_::MultiplicandTraits::kKstrided ? IteratorAdvance::kH
167 : IteratorAdvance::kW,
168 MemorySpace::kGlobal,
169 Index_> {
171 typedef GemmGlobalIteratorAb<TileTraits_, Index_> This_;
172 typedef TileLoadIterator<TileTraits_,
173 typename TileTraits_::Scalar,
174 TileTraits_::MultiplicandTraits::kKstrided ? IteratorAdvance::kH
175 : IteratorAdvance::kW,
176 MemorySpace::kGlobal,
177 Index_>
178 Base;
180 static MatrixLayout::Kind const kLayout = TileTraits_::kLayout;
182 typedef typename TileTraits_::Tile Tile;
184 typedef typename Base::Fragment Fragment;
186 typedef typename TileTraits_::Scalar Scalar;
188 typedef typename TileTraits_::Threads Threads;
190 typedef Index_ Index;
192 typedef long long LongIndex;
194 typedef typename TileTraits_::ThreadOffset ThreadOffset;
196 static IteratorAdvance::Kind const kAdvance = Base::kAdvance;
197
198 typedef cutlass::PredicateVector<ShapeCount<typename Base::Iterations>::kCount> PredicateVector;
199
201 typedef typename Base::Params BaseParams;
202
203 struct Params : public BaseParams {
205 CUTLASS_HOST_DEVICE int initialize(Scalar const* ptr,
206 Index stride_d,
207 Index stride_h) {
208 return BaseParams::initialize(ptr, stride_d, stride_h, kAdvance == IteratorAdvance::kH ? 0 : 1);
209 }
210 };
211
213 Coord<4> thread_offset;
215 Params params;
217 PredicateVector predicates;
218
219 CUTLASS_HOST_DEVICE void initialize_predicates(const Coord<3>& bounds, const Coord<3>& block_offset) {
220 // Setup the masks to control loads.
221 predicates.fill(0);
222
223 // Fill in the bits of the predicate vector.
224 for (int d = 0; d < Base::Iterations::kD; ++d) {
225 for (int h = 0; h < Base::Iterations::kH; ++h) {
226 for (int w = 0; w < Base::Iterations::kW; ++w) {
227 for (int c = 0; c < Base::Iterations::kC; ++c) {
228 bool flag = w * Base::Delta::kW + thread_offset[2] + block_offset[2] < bounds[2];
229 if (kAdvance == IteratorAdvance::kH) {
230 flag =
231 flag &&
232 (h * Base::Delta::kH + d * Base::Delta::kD) + thread_offset[1] + block_offset[1] <
233 bounds[1];
234 } else {
235 flag = flag && (h * Base::Delta::kH) + thread_offset[1] + block_offset[1] < bounds[1];
236 }
237 int const bit = ComputeOffsetFromShape<typename Base::Iterations>::get(d, h, w, c);
238 predicates.set(bit, flag);
239 }
240 }
241 }
242 }
243 }
244
246 CUTLASS_HOST_DEVICE GemmGlobalIteratorAb(Params const& _params,
247 const Coord<3>& threadblock_offset,
248 ThreadOffset thread_offset_func = ThreadOffset())
249 : params(_params) {
250 thread_offset = thread_offset_func();
251 // Setup the pointer.
252 params.pointer += ((threadblock_offset[1] + thread_offset[1]) * params.stride_h +
253 (threadblock_offset[2] + thread_offset[2]));
254
255 }
256
258 CUTLASS_HOST_DEVICE void inc_w() { Base::inc_w(); }
260 CUTLASS_HOST_DEVICE void inc_h() { params.pointer += params.inc_h; }
262 CUTLASS_HOST_DEVICE void inc_d() { params.pointer += params.inc_d; }
264 CUTLASS_HOST_DEVICE void inc_advance() { params.pointer += params.inc_advance; }
265
267 CUTLASS_HOST_DEVICE void load_element(
268 typename Base::AccessType& value, int d, int h, int w, int c) const {
269 int const offset =
270 ComputeOffsetFromStrides<typename Base::ImmediateOffsetStrides>::get(0, 0, w, c);
271 Load<Scalar,
272 Base::kAccessSize,
273 Base::kMemorySpace,
274 Base::kFragmentElementType,
275 typename Base::FragmentElement,
276 Base::Tile::kW,
277 Base::kAccessSize * sizeof(Scalar)>::load(value, params.pointer, offset);
278 }
279
281 CUTLASS_HOST_DEVICE void residue(Index k) {
282 // Update the predicate vector.
283 for (int d = 0; d < Base::Iterations::kD; ++d) {
284 for (int h = 0; h < Base::Iterations::kH; ++h) {
285 for (int w = 0; w < Base::Iterations::kW; ++w) {
286 for (int c = 0; c < Base::Iterations::kC; ++c) {
287 Index offset = 0;
288 if (kAdvance == IteratorAdvance::kH) {
289 offset += thread_offset[1] + h * Base::Delta::kH + d * Base::Delta::kD;
290 } else {
291 offset += thread_offset[2] + w * Base::Delta::kW;
292 }
293
294 int const bit = ComputeOffsetFromShape<typename Base::Iterations>::get(d, h, w, c);
295 if (offset >= k) {
296 predicates.set(bit, false);
297 }
298 }
299 }
300 }
301 }
302 }
303
305 CUTLASS_HOST_DEVICE bool valid(int d, int h, int w, int c) const {
306 int const bit = ComputeOffsetFromShape<typename Base::Iterations>::get(d, h, w, c);
307 return predicates[bit];
308 }
309
311 CUTLASS_HOST_DEVICE GemmGlobalIteratorAb & operator+=(Coord<3> const &offset) {
312
313 LongIndex _offset = offset.template dot<LongIndex>(
314 make_Coord(params.stride_d, params.stride_h, params.stride_w)
315 );
316
317 params.pointer += _offset;
318 return *this;
319 }
320
321 CUTLASS_HOST_DEVICE void add_pointer_offset(Index offset) { params.pointer += offset; }
322
323 CUTLASS_HOST_DEVICE Index stride_advance(void) {
324 Index stride = params.stride_h;
325 if (kAdvance == IteratorAdvance::kW) {
326 stride = params.stride_w;
327 }
328 return stride;
329 }
330
331 template <typename Fragment>
332 CUTLASS_HOST_DEVICE void load_post_increment(Fragment& fragment) {
333 typename Base::FragmentIterator frag_iterator(fragment);
334 for (int d = 0; d < Base::Iterations::kD; ++d) {
335 for (int h = 0; h < Base::Iterations::kH; ++h) {
336 for (int w = 0; w < Base::Iterations::kW; ++w) {
337 for (int c = 0; c < Base::Iterations::kC; ++c) {
338 if (valid(d, h, w, c)) {
339 load_element(
340 reinterpret_cast<typename Base::AccessType&>(frag_iterator.at(d, h, w, c)),
341 d,
342 h,
343 w,
344 c);
345 }
346 }
347 if (w < Base::Iterations::kW - 1) {
348 inc_w();
349 }
350 }
351 if (h < Base::Iterations::kH - 1) {
352 inc_h();
353 }
354 }
355 if (d < Base::Iterations::kD - 1) {
356 inc_d();
357 }
358 }
359 inc_advance();
360 }
361 };
362
364
365 template <typename TileTraits_, typename Index_ = int>
366 struct GemmGlobalIteratorCd : public TileIteratorBase<TileTraits_,
367 typename TileTraits_::Scalar,
368 IteratorAdvance::kH,
369 MemorySpace::kGlobal,
370 Index_> {
372 typedef GemmGlobalIteratorCd<TileTraits_, Index_> This_;
374 typedef TileIteratorBase<TileTraits_,
375 typename TileTraits_::Scalar,
376 IteratorAdvance::kH,
377 MemorySpace::kGlobal,
378 Index_>
379 Base;
380
382 static MatrixLayout::Kind const kLayout = TileTraits_::kLayout;
383
385 typedef typename TileTraits_::Scalar Scalar;
387 typedef typename TileTraits_::Pointer Pointer;
389 typedef typename TileTraits_::Threads Threads;
391 typedef Index_ Index;
393 typedef long long LongIndex;
395 typedef typename TileTraits_::ThreadOffset ThreadOffset;
396
398 struct Params {
400 Pointer pointer;
402 long long stride_d;
404 Index stride_h;
406 Index inc_advance, inc_h;
408 Index predicate_inc_advance, predicate_inc_h;
410 Index predicate_offset;
411
413 CUTLASS_HOST_DEVICE int initialize(Pointer pointer,
414 int stride_d_,
415 Index ldm,
416 Index bound,
417 Index epilogue_stride_w,
418 Index epilogue_delta_w) {
419 // The pointer.
420 this->pointer = pointer;
421 // Stride per batch
422 stride_d = stride_d_;
423 // Each column of the matrix.
424 stride_h = TileTraits_::ThreadsDelta::kH * ldm;
425 // Each thread output 1 column per iteration. The stride between columns is given by the
426 // number of scalars that are loaded per LDS for B.
427 inc_h = ldm * TileTraits_::kStrideH;
428 inc_advance =
429 (ldm - ldm * TileTraits_::kStrideH * (Base::Iterations::kH - 1)) + epilogue_stride_w;
430
431 predicate_offset = bound;
432 predicate_inc_h = TileTraits_::kStrideH;
433 predicate_inc_advance =
434 -((TileTraits_::kStrideH * (Base::Iterations::kH - 1) - 1) + epilogue_delta_w);
435
436 return 0;
437 }
438
439 CUTLASS_HOST_DEVICE int initialize(Pointer pointer, long long _stride_d, Index _stride_h,
440 Index _inc_advance, Index _inc_h, Index _predicate_inc_advance, Index _predicate_inc_h,
441 Index _predicate_offset) {
442 this->pointer = pointer;
443 stride_d = _stride_d;
444 stride_h = _stride_h;
445 inc_advance = _inc_advance;
446 inc_h = _inc_h;
447 predicate_inc_advance = _predicate_inc_advance;
448 predicate_inc_h = _predicate_inc_h;
449 predicate_offset = _predicate_offset;
450
451 return 0;
452 }
453 };
454
456 Params params;
458 Coord<4> thread_offset;
460 cutlass::PredicateVector<Base::Iterations::kW> predicates;
461
463 CUTLASS_HOST_DEVICE GemmGlobalIteratorCd(Params const& _params,
464 const Coord<3>& bounds,
465 const Coord<3>& block,
466 int offset = 0,
467 int pred_offset = 0,
468 ThreadOffset thread_offset_func = ThreadOffset())
469 : params(_params) {
470 thread_offset = thread_offset_func();
471 // Each warp works on a different column of the tile.
472 int const h = thread_offset[1] + block[1];
473 // Each lane writes a different element.
474 int const w = thread_offset[2] + block[2];
475 // Setup the pointer.
476 params.pointer += ((h * params.stride_h + w) + offset);
477
478 // Prepare the vector of predicates.
479 for (int i = 0; i < Base::Iterations::kW; ++i) {
480 predicates.set(i, w + i * Base::Delta::kW < bounds[2]);
481 }
482 params.predicate_offset -= (h + pred_offset);
483 }
484
486 CUTLASS_HOST_DEVICE void inc_c() {}
488 CUTLASS_HOST_DEVICE void inc_w() {}
490 CUTLASS_HOST_DEVICE void inc_h() {
491 params.pointer += params.inc_h;
492 params.predicate_offset -= params.predicate_inc_h;
493 }
495 CUTLASS_HOST_DEVICE void inc_d() {}
497 CUTLASS_HOST_DEVICE void inc_advance() {
498 params.pointer += params.inc_advance;
499 params.predicate_offset -= params.predicate_inc_advance;
500 }
501
503 CUTLASS_HOST_DEVICE GemmGlobalIteratorCd & operator+=(Coord<3> const &offset) {
504 LongIndex _offset = offset.template dot<LongIndex>(
505 make_Coord(params.stride_d, params.stride_h, 1)
506 );
507 params.pointer += _offset;
508 return *this;
509 }
510
512 CUTLASS_HOST_DEVICE void load_element(
513 typename Base::AccessType& value, int d, int h, int w, int c) const {
514 int const offset =
515 ComputeOffsetFromStrides<typename Base::ImmediateOffsetStrides>::get(d, h, w, c);
516 Load<Scalar,
517 Base::kAccessSize,
518 Base::kMemorySpace,
519 Base::kFragmentElementType,
520 typename Base::FragmentElement,
521 Base::Tile::kW,
522 Base::kAccessSize * sizeof(Scalar)>::load(value, params.pointer, offset);
523 }
524
526 CUTLASS_HOST_DEVICE void store_element(
527 typename Base::AccessType const& value, int d, int h, int w, int c) {
528 int const offset =
529 ComputeOffsetFromStrides<typename Base::ImmediateOffsetStrides>::get(d, h, w, c);
530 Store<Scalar,
531 Base::kAccessSize,
532 Base::kMemorySpace,
533 Base::kFragmentElementType,
534 typename Base::FragmentElement,
535 Base::Tile::kW,
536 Base::kAccessSize * sizeof(Scalar)>::store(value, params.pointer, offset);
537 }
538
540 CUTLASS_HOST_DEVICE bool valid(int d, int h, int w, int c) const {
541 return predicates.at(w) && params.predicate_offset > 0;
542 }
543
545 CUTLASS_HOST_DEVICE void add_pointer_offset(LongIndex offset) { params.pointer += offset; }
546
548 template <typename Fragment>
549 CUTLASS_HOST_DEVICE void load_post_increment(Fragment& fragment) {
550 typename Base::FragmentIterator frag_iterator(fragment);
551 for (int d = 0; d < Base::Iterations::kD; ++d) {
552 for (int h = 0; h < Base::Iterations::kH; ++h) {
553 for (int w = 0; w < Base::Iterations::kW; ++w) {
554 for (int c = 0; c < Base::Iterations::kC; ++c) {
555 if (valid(d, h, w, c)) {
556 load_element(
557 reinterpret_cast<typename Base::AccessType&>(frag_iterator.at(d, h, w, c)),
558 d,
559 h,
560 w,
561 c);
562 }
563 }
564 if (w < Base::Iterations::kW - 1) {
565 inc_w();
566 }
567 }
568 if (h < Base::Iterations::kH - 1) {
569 inc_h();
570 }
571 }
572 if (d < Base::Iterations::kD - 1) {
573 inc_d();
574 }
575 }
576 inc_advance();
577 }
578
579 template <typename Fragment>
580 CUTLASS_HOST_DEVICE void store_post_increment(Fragment& fragment) {
581 typename Base::FragmentIterator frag_iterator(fragment);
582 for (int d = 0; d < Base::Iterations::kD; ++d) {
583 for (int h = 0; h < Base::Iterations::kH; ++h) {
584 for (int w = 0; w < Base::Iterations::kW; ++w) {
585 for (int c = 0; c < Base::Iterations::kC; ++c) {
586 if (valid(d, h, w, c)) {
587 store_element(
588 reinterpret_cast<typename Base::AccessType&>(frag_iterator.at(d, h, w, c)),
589 d,
590 h,
591 w,
592 c);
593 }
594 }
595 if (w < Base::Iterations::kW - 1) {
596 inc_w();
597 }
598 }
599 if (h < Base::Iterations::kH - 1) {
600 inc_h();
601 }
602 }
603 if (d < Base::Iterations::kD - 1) {
604 inc_d();
605 }
606 }
607 inc_advance();
608 }
609 };
610
612
613 } // namespace gemm
614 } // namespace cutlass
cutlass::gemm::GemmGlobalTileCdTraits
Definition: gemm_global_tile.h:120
cutlass::gemm::GemmGlobalTileTraits::Delta
Shape< 0, Threads::kH, Threads::kW *kAccessSize > Delta
The strides in each dimension between different loads/stores.
Definition: gemm_global_tile.h:92
cutlass::gemm::GemmGlobalIteratorCd::operator+=
CUTLASS_HOST_DEVICE GemmGlobalIteratorCd & operator+=(Coord< 3 > const &offset)
Adds a vector offset to the iterator.
Definition: gemm_global_tile.h:503
cutlass::gemm::GemmGlobalIteratorCd::add_pointer_offset
CUTLASS_HOST_DEVICE void add_pointer_offset(LongIndex offset)
add pointer offset
Definition: gemm_global_tile.h:545
cutlass::gemm::GemmGlobalIteratorCd::Params::initialize
CUTLASS_HOST_DEVICE int initialize(Pointer pointer, long long _stride_d, Index _stride_h, Index _inc_advance, Index _inc_h, Index _predicate_inc_advance, Index _predicate_inc_h, Index _predicate_offset)
Definition: gemm_global_tile.h:439
cutlass::gemm::GemmGlobalIteratorAb::load_post_increment
CUTLASS_HOST_DEVICE void load_post_increment(Fragment &fragment)
Definition: gemm_global_tile.h:332
cutlass::gemm::GemmGlobalIteratorCd::Params::inc_advance
Index inc_advance
The strides to increment the pointer.
Definition: gemm_global_tile.h:406
cutlass
Definition: convert.h:33
cutlass::gemm::GemmGlobalIteratorAb::PredicateVector
cutlass::PredicateVector< ShapeCount< typename Base::Iterations >::kCount > PredicateVector
Definition: gemm_global_tile.h:198
cutlass::gemm::GemmGlobalIteratorAb::kLayout
static MatrixLayout::Kind const kLayout
The layout.
Definition: gemm_global_tile.h:180
cutlass::platform::remove_const::type
T type
Definition: platform.h:377
cutlass::TileLoadIterator::kAccessSize
static int const kAccessSize
The number of scalars accessed per load/store.
Definition: tile_iterator.h:467
cutlass::gemm::GemmGlobalIteratorAb::BaseParams
Base::Params BaseParams
Iterator parameters type.
Definition: gemm_global_tile.h:201
cutlass::gemm::GemmGlobalIteratorCd::inc_c
CUTLASS_HOST_DEVICE void inc_c()
Increment the pointer in the C dimension.
Definition: gemm_global_tile.h:486
cutlass::gemm::GemmGlobalIteratorCd::Index
Index_ Index
The index.
Definition: gemm_global_tile.h:391
tile_iterator.h
Defines the Tile Traits concept and iterators for loading and storing to tiles efficiently.
cutlass::gemm::GemmGlobalTileTraits::VectorizedTile
ReshapeTile< Tile_, kAccessSize_ >::Tile VectorizedTile
The vectorized tile shape.
Definition: gemm_global_tile.h:86
cutlass::gemm::GemmGlobalIteratorCd::This_
GemmGlobalIteratorCd< TileTraits_, Index_ > This_
This class.
Definition: gemm_global_tile.h:372
cutlass::gemm::GemmGlobalIteratorCd::kLayout
static MatrixLayout::Kind const kLayout
The layout.
Definition: gemm_global_tile.h:382
cutlass::gemm::GemmGlobalTileTraits
Definition: gemm_global_tile.h:70
cutlass::gemm::GemmGlobalTileTraits::Pointer
Scalar_ * Pointer
The pointer.
Definition: gemm_global_tile.h:78
cutlass::TileIteratorBase::FragmentIterator
FragmentIterator< Fragment, Iterations, AccessType > FragmentIterator
The fragment iterator.
Definition: tile_iterator.h:202
coord.h
A Coord is a coordinate of arbitrary rank into a tensor or matrix.
cutlass::IteratorAdvance::Kind
Kind
Definition: tile_iterator.h:65
cutlass::gemm::GemmGlobalIteratorAb::Params::initialize
CUTLASS_HOST_DEVICE int initialize(Scalar const *ptr, Index stride_d, Index stride_h)
Initializes params to load a strip-mined tile, given pointer and stride_h.
Definition: gemm_global_tile.h:205
cutlass::PredicateVector::at
CUTLASS_HOST_DEVICE bool at(int idx) const
Accesses a bit within the predicate vector.
Definition: predicate_vector.h:357
cutlass::MemorySpace::kGlobal
Definition: load_store.h:42
cutlass::make_Coord
CUTLASS_HOST_DEVICE Coord< 1 > make_Coord(int _0)
Helper to make a 2-element coordinate.
Definition: coord.h:368
cutlass::gemm::GemmGlobalTileTraits::ThreadsDelta
Shape< 1, 1, VectorizedTile::kC > ThreadsDelta
The relative offset between two elements in the H/W dimension in adjacent threads.
Definition: gemm_global_tile.h:90
cutlass::gemm::GemmGlobalTileTraits::MultiplicandTraits
GemmMultiplicandTraits< Tile, kOperand, kLayout > MultiplicandTraits
Definition: gemm_global_tile.h:103
cutlass::TileLoadIterator::FragmentElement
FragmentElement_ FragmentElement
Fragment element.
Definition: tile_iterator.h:428
cutlass::gemm::GemmGlobalTileTraits::kMemorySpace
static MemorySpace::Kind const kMemorySpace
The memory space.
Definition: gemm_global_tile.h:82
cutlass::gemm::GemmGlobalIteratorAb::LongIndex
long long LongIndex
Long index.
Definition: gemm_global_tile.h:192
cutlass::gemm::GemmGlobalIteratorAb::load_element
CUTLASS_HOST_DEVICE void load_element(typename Base::AccessType &value, int d, int h, int w, int c) const
Loads a single fragment element from memory.
Definition: gemm_global_tile.h:267
cutlass::gemm::GemmGlobalIteratorCd::Base
TileIteratorBase< TileTraits_, typename TileTraits_::Scalar, IteratorAdvance::kH, MemorySpace::kGlobal, Index_ > Base
The base class.
Definition: gemm_global_tile.h:379
cutlass::TileIteratorBase< TileTraits_, TileTraits_::Scalar, IteratorAdvance::kH, MemorySpace::kGlobal, Index_ >::Fragment
Fragment< FragmentElement, ShapeCount< Iterations >::kCount *kAccessSize > Fragment
The fragment.
Definition: tile_iterator.h:199
cutlass::MatrixLayout::Kind
Kind
Enumeration defining fundamental contiguous layouts.
Definition: matrix_traits.h:159
cutlass::gemm::GemmGlobalTileCdTraits::Delta
Shape< 0, 0, Base::Delta::kW, Base::Delta::kC > Delta
Override the strides in each dimension between different loads/stores.
Definition: gemm_global_tile.h:138
cutlass::gemm::GemmGlobalIteratorCd::Params::predicate_inc_h
Index predicate_inc_h
Definition: gemm_global_tile.h:408
cutlass::TileLoadIterator::Params::initialize
CUTLASS_HOST_DEVICE int initialize()
Definition: tile_iterator.h:590
cutlass::TileLoadIterator::kMemorySpace
static MemorySpace::Kind const kMemorySpace
Source or destination memory space.
Definition: tile_iterator.h:437
cutlass::gemm::GemmGlobalTileTraits::Tile
Tile_ Tile
The tile shape.
Definition: gemm_global_tile.h:84
cutlass::TileLoadIterator::Fragment
Base::Fragment Fragment
Fragment definition.
Definition: tile_iterator.h:470
cutlass::IteratorAdvance::kH
Definition: tile_iterator.h:65
cutlass::gemm::GemmGlobalIteratorCd::load_post_increment
CUTLASS_HOST_DEVICE void load_post_increment(Fragment &fragment)
Loads and increments iterator.
Definition: gemm_global_tile.h:549
cutlass::gemm::GemmGlobalIteratorAb::Base
TileLoadIterator< TileTraits_, typename TileTraits_::Scalar, TileTraits_::MultiplicandTraits::kKstrided ? IteratorAdvance::kH :IteratorAdvance::kW, MemorySpace::kGlobal, Index_ > Base
The base class.
Definition: gemm_global_tile.h:178
cutlass::gemm::GemmGlobalIteratorAb::Params
Definition: gemm_global_tile.h:203
cutlass::TileIteratorBase::Params::inc_d
Index inc_d
Definition: tile_iterator.h:226
cutlass::gemm::GemmGlobalIteratorCd::GemmGlobalIteratorCd
CUTLASS_HOST_DEVICE GemmGlobalIteratorCd(Params const &_params, const Coord< 3 > &bounds, const Coord< 3 > &block, int offset=0, int pred_offset=0, ThreadOffset thread_offset_func=ThreadOffset())
Ctor.
Definition: gemm_global_tile.h:463
cutlass::GemmOperand::kC
Definition: matrix_traits.h:357
cutlass::ComputeOffsetFromShape::get
static CUTLASS_HOST_DEVICE int get(int d, int h, int w, int c)
Definition: shape.h:181
platform.h
C++ features that may be otherwise unimplemented for CUDA device functions.
cutlass::gemm::GemmGlobalIteratorAb
Definition: gemm_global_tile.h:163
cutlass::gemm::GemmGlobalTileCdTraits::Base
GemmGlobalTileTraits< GemmOperand::kC, MatrixLayout::kColumnMajor, Scalar_, Tile_, Threads_, kAccessSize_ > Base
The base class.
Definition: gemm_global_tile.h:133
cutlass::MemorySpace::Kind
Kind
Definition: load_store.h:39
cutlass::gemm::GemmGlobalIteratorAb::inc_d
CUTLASS_HOST_DEVICE void inc_d()
Increment the pointer in the D dimension.
Definition: gemm_global_tile.h:262
cutlass::gemm::GemmGlobalIteratorAb::kAdvance
static IteratorAdvance::Kind const kAdvance
Specifies in which dimension post-increment accesses advance.
Definition: gemm_global_tile.h:196
cutlass::gemm::GemmGlobalIteratorAb::Threads
TileTraits_::Threads Threads
The threads.
Definition: gemm_global_tile.h:188
cutlass::gemm::GemmGlobalTileCdTraits::kStrideH
static int const kStrideH
The stride in the H dimension.
Definition: gemm_global_tile.h:136
cutlass::Shape::kH
static int const kH
The height of the cube.
Definition: shape.h:68
cutlass::Store
Definition: load_store.h:178
cutlass::gemm::ReshapeThreads< Tile_, Threads_, true >::Threads
Shape< Threads_::kD, Threads_::kH *Threads_::kW/Tile_::kW, Tile_::kW, 1 > Threads
Definition: gemm_global_tile.h:59
cutlass::gemm::GemmGlobalIteratorCd::Params::predicate_inc_advance
Index predicate_inc_advance
The strides to increment the predicate offset.
Definition: gemm_global_tile.h:408
cutlass::gemm::GemmGlobalTileTraits::kOperand
static GemmOperand::Kind const kOperand
Identity of the operand.
Definition: gemm_global_tile.h:72
cutlass::TileIteratorBase::Params::stride_h
Index stride_h
Definition: tile_iterator.h:223
predicate_vector.h
Defines container classes and iterators for managing a statically sized vector of boolean predicates...
cutlass::gemm::GemmGlobalIteratorCd::valid
CUTLASS_HOST_DEVICE bool valid(int d, int h, int w, int c) const
Test the validity of the.
Definition: gemm_global_tile.h:540
cutlass::TileIteratorBase::kFragmentElementType
static FragmentElementType::Kind const kFragmentElementType
Specifies iterator storage fragment type (Scalar or WmmaMatrix)
Definition: tile_iterator.h:158
cutlass::gemm::GemmGlobalIteratorCd::inc_d
CUTLASS_HOST_DEVICE void inc_d()
Increment the pointer in the D dimension.
Definition: gemm_global_tile.h:495
cutlass::TileLoadIterator
An iterator implementing Tile Load Iterator Concept for loading a tile from memory.
Definition: tile_iterator.h:402
cutlass::gemm::ReshapeThreads< Tile_, Threads_, true >
Definition: gemm_global_tile.h:58
cutlass::gemm::GemmGlobalIteratorAb::GemmGlobalIteratorAb
CUTLASS_HOST_DEVICE GemmGlobalIteratorAb(Params const &_params, const Coord< 3 > &threadblock_offset, ThreadOffset thread_offset_func=ThreadOffset())
Ctor.
Definition: gemm_global_tile.h:246
cutlass::gemm::GemmGlobalIteratorAb::predicates
PredicateVector predicates
The predicates.
Definition: gemm_global_tile.h:217
cutlass::gemm::GemmGlobalTileTraits::Scalar
platform::remove_const< Scalar_ >::type Scalar
The scalar.
Definition: gemm_global_tile.h:76
reshape_tile.h
Defines a type for restructuring a tile.
gemm_operand.h
Defines constant expressions for mapping GEMM problem size and strides onto pitch-linear memory...
cutlass::gemm::GemmGlobalIteratorAb::inc_h
CUTLASS_HOST_DEVICE void inc_h()
Increment the pointer in the H dimension.
Definition: gemm_global_tile.h:260
cutlass::gemm::GemmGlobalIteratorAb::operator+=
CUTLASS_HOST_DEVICE GemmGlobalIteratorAb & operator+=(Coord< 3 > const &offset)
Adds a vector offset to the iterator.
Definition: gemm_global_tile.h:311
cutlass::gemm::GemmGlobalIteratorAb::Fragment
Base::Fragment Fragment
Fragment type loaded by the iterator.
Definition: gemm_global_tile.h:184
cutlass::gemm::GemmGlobalIteratorCd::Threads
TileTraits_::Threads Threads
The threads.
Definition: gemm_global_tile.h:389
cutlass::gemm::GemmGlobalTileCdTraits::ThreadOffset::operator()
CUTLASS_HOST_DEVICE Coord< 4 > operator()() const
Definition: gemm_global_tile.h:151
cutlass::gemm::GemmMultiplicandTraits
Definition: gemm_operand.h:67
cutlass::gemm::GemmGlobalTileTraits::ThreadOffset
Computes the thread offset in (H, W) based on thread ID.
Definition: gemm_global_tile.h:106
cutlass::gemm::GemmGlobalIteratorAb::inc_w
CUTLASS_HOST_DEVICE void inc_w()
Increment the pointer in the W dimension.
Definition: gemm_global_tile.h:258
cutlass::TileLoadIterator< TileTraits_, TileTraits_::Scalar, TileTraits_::MultiplicandTraits::kKstrided ? IteratorAdvance::kH :IteratorAdvance::kW, MemorySpace::kGlobal, Index_ >::load
CUTLASS_HOST_DEVICE void load(Fragment &fragment, PredicateIterator pred_it) const
Loads a fragment without advancing the iterator..
Definition: tile_iterator.h:776
cutlass::PredicateVector::fill
CUTLASS_HOST_DEVICE void fill(bool value=true)
Fills all predicates with a given value.
Definition: predicate_vector.h:344
cutlass::gemm::GemmGlobalIteratorCd::store_post_increment
CUTLASS_HOST_DEVICE void store_post_increment(Fragment &fragment)
Definition: gemm_global_tile.h:580
cutlass::gemm::GemmGlobalTileCdTraits::Threads
Base::Threads Threads
Definition: gemm_global_tile.h:142
cutlass::gemm::GemmGlobalIteratorCd::Params::stride_h
Index stride_h
The stride in the H dimension to setup the thread in the block.
Definition: gemm_global_tile.h:404
CUTLASS_HOST_DEVICE
#define CUTLASS_HOST_DEVICE
Definition: cutlass.h:46
cutlass::gemm::GemmGlobalIteratorAb::inc_advance
CUTLASS_HOST_DEVICE void inc_advance()
Increment the pointer to move to the next iteration.
Definition: gemm_global_tile.h:264
cutlass::gemm::GemmGlobalTileTraits::ThreadOffset::operator()
CUTLASS_HOST_DEVICE Coord< 4 > operator()() const
Definition: gemm_global_tile.h:108
cutlass::gemm::GemmGlobalIteratorCd::load_element
CUTLASS_HOST_DEVICE void load_element(typename Base::AccessType &value, int d, int h, int w, int c) const
Loads a single fragment element from memory.
Definition: gemm_global_tile.h:512
cutlass::TileIteratorBase::Params::inc_h
Index inc_h
Definition: tile_iterator.h:227
cutlass::TileIteratorBase::Params::stride_d
Index stride_d
Definition: tile_iterator.h:222
cutlass::gemm::GemmGlobalTileTraits::ImmediateOffsetStrides
Shape< 0, 0, Threads::kW *ThreadsDelta::kW, kAccessSize > ImmediateOffsetStrides
Strides for immediate offset computation.
Definition: gemm_global_tile.h:95
cutlass::PredicateVector
Statically sized array of bits implementing.
Definition: predicate_vector.h:105
cutlass::Vector
Definition: vector.h:62
cutlass::Load
Definition: load_store.h:60
cutlass::gemm::GemmGlobalIteratorCd::ThreadOffset
TileTraits_::ThreadOffset ThreadOffset
The thread offset.
Definition: gemm_global_tile.h:395
cutlass::Shape
A Shape implementing Layout Concept describing the dimensions of a cube.
Definition: shape.h:64
cutlass::MatrixLayout::kColumnMajor
Definition: matrix_traits.h:159
cutlass::gemm::GemmGlobalTileCdTraits::ImmediateOffsetStrides
Base::ImmediateOffsetStrides ImmediateOffsetStrides
Definition: gemm_global_tile.h:146
cutlass::gemm::GemmGlobalIteratorCd::LongIndex
long long LongIndex
The index.
Definition: gemm_global_tile.h:393
cutlass::gemm::GemmGlobalIteratorCd::Scalar
TileTraits_::Scalar Scalar
The scalar.
Definition: gemm_global_tile.h:385
cutlass::gemm::GemmGlobalIteratorCd::Params::inc_h
Index inc_h
Definition: gemm_global_tile.h:406
cutlass::gemm::GemmGlobalIteratorCd::predicates
cutlass::PredicateVector< Base::Iterations::kW > predicates
The predicates for the row.
Definition: gemm_global_tile.h:460
cutlass::TileIteratorBase::kMemorySpace
static MemorySpace::Kind const kMemorySpace
Source or destination memory space.
Definition: tile_iterator.h:161
cutlass::gemm::GemmGlobalIteratorCd::Params::pointer
Pointer pointer
The pointer.
Definition: gemm_global_tile.h:400
cutlass::gemm::GemmGlobalIteratorAb::This_
GemmGlobalIteratorAb< TileTraits_, Index_ > This_
This class.
Definition: gemm_global_tile.h:171
cutlass::Coord< 4 >
cutlass::gemm::GemmGlobalIteratorAb::Tile
TileTraits_::Tile Tile
The tile.
Definition: gemm_global_tile.h:182
cutlass::gemm::GemmGlobalTileTraits::Threads
ReshapeThreads< VectorizedTile, Threads_ >::Threads Threads
The threads shape.
Definition: gemm_global_tile.h:88
cutlass::gemm::GemmGlobalTileTraits::Iterations
Shape< 1, VectorizedTile::kH/Threads::kH, VectorizedTile::kW/Threads::kW, VectorizedTile::kC/kAccessSize > Iterations
The number of iterations needed to load/store the tile.
Definition: gemm_global_tile.h:101
cutlass::gemm::GemmGlobalIteratorCd::inc_advance
CUTLASS_HOST_DEVICE void inc_advance()
Increment the pointer to move to the next iteration.
Definition: gemm_global_tile.h:497
cutlass::gemm::GemmGlobalIteratorCd::inc_h
CUTLASS_HOST_DEVICE void inc_h()
Increment the pointer in the H dimension.
Definition: gemm_global_tile.h:490
cutlass::gemm::GemmGlobalIteratorCd::inc_w
CUTLASS_HOST_DEVICE void inc_w()
Increment the pointer in the W dimension.
Definition: gemm_global_tile.h:488
cutlass::TileIteratorBase
Iterator for accessing a stripmined tile in memory.
Definition: tile_iterator.h:144
cutlass::TileLoadIterator::Params::pointer
Scalar const * pointer
Pointer to memory.
Definition: tile_iterator.h:499
cutlass::gemm::GemmGlobalIteratorCd::params
Params params
Parameters.
Definition: gemm_global_tile.h:456
cutlass::gemm::GemmGlobalIteratorCd
Definition: gemm_global_tile.h:366
cutlass::TileLoadIterator::kAdvance
static IteratorAdvance::Kind const kAdvance
Specifies in which dimension post-increment accesses advance.
Definition: tile_iterator.h:431
cutlass::gemm::GemmGlobalIteratorCd::thread_offset
Coord< 4 > thread_offset
Offset of an individual lane from the start of the tile.
Definition: gemm_global_tile.h:458
cutlass::TileIteratorBase::FragmentElement
FragmentElement_ FragmentElement
Fragment element.
Definition: tile_iterator.h:152
cutlass::gemm::GemmGlobalIteratorAb::ThreadOffset
TileTraits_::ThreadOffset ThreadOffset
The thread offset.
Definition: gemm_global_tile.h:194
cutlass::Shape::kW
static int const kW
The width of the cube.
Definition: shape.h:70
cutlass::PredicateVector::set
CUTLASS_HOST_DEVICE void set(int idx, bool value=true)
Set a bit within the predicate vector.
Definition: predicate_vector.h:365
cutlass::gemm::GemmGlobalIteratorAb::add_pointer_offset
CUTLASS_HOST_DEVICE void add_pointer_offset(Index offset)
Definition: gemm_global_tile.h:321
cutlass::TileIteratorBase::kAccessSize
static int const kAccessSize
The number of scalars accessed per load/store.
Definition: tile_iterator.h:188
cutlass::TileLoadIterator::Params
Parameters.
Definition: tile_iterator.h:497
cutlass::ComputeOffsetFromStrides::get
static CUTLASS_HOST_DEVICE int get(int d, int h, int w, int c)
Definition: shape.h:199
cutlass::gemm::GemmGlobalTileCdTraits::ThreadOffset
Computes the thread offset in (H, W) based on thread ID.
Definition: gemm_global_tile.h:149
cutlass::gemm::GemmGlobalIteratorCd::Params::stride_d
long long stride_d
The stride in the D dimension.
Definition: gemm_global_tile.h:402
cutlass::gemm::GemmGlobalIteratorAb::stride_advance
CUTLASS_HOST_DEVICE Index stride_advance(void)
Definition: gemm_global_tile.h:323
cutlass::TileLoadIterator::inc_w
CUTLASS_HOST_DEVICE void inc_w()
Increment in the W dimension.
Definition: tile_iterator.h:686
cutlass::gemm::GemmGlobalTileTraits::kAccessSize
static int const kAccessSize
The number of scalars per LDG/STG.
Definition: gemm_global_tile.h:80
cutlass::ReshapeTile::Tile
Tile_ Tile
Definition: reshape_tile.h:43
cutlass::IteratorAdvance::kW
Definition: tile_iterator.h:65
cutlass::gemm::GemmGlobalTileCdTraits::Iterations
Base::Iterations Iterations
Definition: gemm_global_tile.h:140
cutlass::gemm::GemmGlobalIteratorAb::Index
Index_ Index
The index.
Definition: gemm_global_tile.h:190
cutlass::gemm::GemmGlobalIteratorCd::Pointer
TileTraits_::Pointer Pointer
The pointer.
Definition: gemm_global_tile.h:387
cutlass::gemm::GemmGlobalIteratorAb::valid
CUTLASS_HOST_DEVICE bool valid(int d, int h, int w, int c) const
Is the valid?
Definition: gemm_global_tile.h:305
cutlass::TileLoadIterator::kFragmentElementType
static FragmentElementType::Kind const kFragmentElementType
Specifies type of iterator fragment storage (Salar or WmmaMatrix)
Definition: tile_iterator.h:434
cutlass::GemmOperand::Kind
Kind
Definition: matrix_traits.h:357
cutlass::gemm::GemmGlobalIteratorAb::Scalar
TileTraits_::Scalar Scalar
The scalar.
Definition: gemm_global_tile.h:186
cutlass::TileIteratorBase::Params::inc_advance
Index inc_advance
Definition: tile_iterator.h:230
cutlass::gemm::ReshapeThreads::Threads
Threads_ Threads
Definition: gemm_global_tile.h:54
cutlass::gemm::GemmGlobalIteratorAb::params
Params params
The parameters.
Definition: gemm_global_tile.h:215
matrix_traits.h
Defines properties of matrices used to denote layout and operands to GEMM kernels.
cutlass::gemm::GemmGlobalIteratorCd::Params::initialize
CUTLASS_HOST_DEVICE int initialize(Pointer pointer, int stride_d_, Index ldm, Index bound, Index epilogue_stride_w, Index epilogue_delta_w)
Setup the params.
Definition: gemm_global_tile.h:413
cutlass::gemm::GemmGlobalIteratorAb::initialize_predicates
CUTLASS_HOST_DEVICE void initialize_predicates(const Coord< 3 > &bounds, const Coord< 3 > &block_offset)
Definition: gemm_global_tile.h:219
cutlass::gemm::GemmGlobalIteratorCd::Params
The params.
Definition: gemm_global_tile.h:398
cutlass::gemm::GemmGlobalTileCdTraits::ThreadsDelta
Base::ThreadsDelta ThreadsDelta
Definition: gemm_global_tile.h:144
cutlass::gemm::GemmGlobalIteratorAb::thread_offset
Coord< 4 > thread_offset
Offset of an individual lane from the start of the tile.
Definition: gemm_global_tile.h:213
cutlass::TileLoadIterator::FragmentIterator
Base::FragmentIterator FragmentIterator
Fragment iterator definition.
Definition: tile_iterator.h:473
cutlass::gemm::GemmGlobalIteratorAb::residue
CUTLASS_HOST_DEVICE void residue(Index k)
That&#39;s the residue! Update the predicates.
Definition: gemm_global_tile.h:281
cutlass::gemm::GemmGlobalIteratorCd::Params::predicate_offset
Index predicate_offset
The column offset to compute the predicate for the columns.
Definition: gemm_global_tile.h:410
cutlass::gemm::GemmGlobalTileTraits::kLayout
static MatrixLayout::Kind const kLayout
The layout.
Definition: gemm_global_tile.h:74
cutlass::gemm::GemmGlobalIteratorCd::store_element
CUTLASS_HOST_DEVICE void store_element(typename Base::AccessType const &value, int d, int h, int w, int c)
Stores a single fragment element into memory.
Definition: gemm_global_tile.h:526
cutlass::TileIteratorBase::Params::stride_w
Index stride_w
Definition: tile_iterator.h:224