Cutlass CUDA Templates for Linear Algebra Subroutines and Solvers

igemm_traits.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 **************************************************************************************************/
30 #pragma once
31
32 #include "cutlass/convert.h"
33 #include "cutlass/gemm/gemm.h"
34 #include "cutlass/gemm/gemm_epilogue.h"
35 #include "cutlass/gemm/gemm_epilogue_traits.h"
36 #include "cutlass/gemm/gemm_global_tile.h"
37 #include "cutlass/gemm/gemm_shared_tile.h"
38 #include "cutlass/gemm/gemm_traits.h"
39 #include "cutlass/gemm/igemm_epilogue.h"
40 #include "cutlass/gemm/igemm_global_tile.h"
41 #include "cutlass/gemm/igemm_multiply_add.h"
42 #include "cutlass/gemm/igemm_swizzle.h"
43 #include "cutlass/reshape_tile.h"
44
45 namespace cutlass {
46 namespace gemm {
47
49
50 template <
52 typename OutputTile_,
54 typename ScalarD_,
56 typename ThreadGemmShape_>
57 struct IgemmConfig : public GemmConfig<
59 int8_t,
61 int8_t,
63 ScalarD_,
65 ScalarD_,
67 OutputTile_,
69 ThreadMultiplyAdd<ThreadGemmShape_, Shape<1, 4, 8>, int8_t, int8_t, int>,
71 4,
73 4,
75 16,
77 4,
79 4,
81 16,
83 1,
85 4,
87 1,
89 2,
91 false,
93 false,
95 false> {};
96
98
99 template <typename OutputTile_, typename ThreadGemmShape_>
100 struct IgemmConfig<OutputTile_, int8_t, ThreadGemmShape_>
101 : public GemmConfig<
103 int8_t,
105 int8_t,
107 int8_t,
109 int8_t,
111 OutputTile_,
113 ThreadMultiplyAdd<ThreadGemmShape_, Shape<1, 4, 8>, int8_t, int8_t, int>,
115 4,
117 4,
119 16,
121 4,
123 4,
125 16,
127 4,
129 4,
131 4,
133 2,
135 false,
137 true,
139 false> {};
140
142
143 template <enum MatrixLayout::Kind kLayout_, typename GemmConfig_, typename Index_>
144 struct IgemmTileTraitsHelperA : public GemmTileTraitsHelperA<kLayout_, GemmConfig_> {};
145
147
148 template <typename GemmConfig_, typename Index_>
149 struct IgemmTileTraitsHelperA<MatrixLayout::kColumnMajor, GemmConfig_, Index_>
150 : public GemmTileTraitsHelperA<MatrixLayout::kColumnMajor, GemmConfig_> {
152 typedef GemmTileTraitsHelperA<MatrixLayout::kColumnMajor, GemmConfig_> Base;
153
155 static int const kScalarsPerStsA = 16;
156
158 typedef IgemmGlobalTileTraits<
159 GemmOperand::kA,
160 // The layout.
161 MatrixLayout::kColumnMajor,
162 // The pointer is float const.
163 int8_t const,
164 // The tile has size KxM in GEMM's terminology.
165 Shape<1, GemmConfig_::OutputTile::kD, GemmConfig_::OutputTile::kW>,
166 // The threads are distributed as warps x 32 (the traits may reorganize).
167 Shape<1, ShapeCount<typename GemmConfig_::Warps>::kCount, GemmConfig_::kWarpSize>,
168 // The number of scalars per LDG (LDG.32 or LDG.128, etc).
169 GemmConfig_::kScalarsPerLdgA>
170 GlobalTileTraits;
171
173 typedef GemmGlobalIteratorAb<GlobalTileTraits, Index_> GlobalLoadIterator;
174
176 typedef GemmSharedStoreTileAbTraits<
177 // The pointer is float.
178 int8_t,
179 // The tile has size KxM in GEMM's terminology.
180 Shape<GemmConfig_::kStages, GemmConfig_::OutputTile::kD / 4, GemmConfig_::OutputTile::kW * 4>,
181 // The threads are distributed as warps x 32 (the traits may reorganize).
182 typename GlobalTileTraits::Threads,
183 // The number of scalars per STS (STS.32 or STS.128, etc).
184 kScalarsPerStsA>
185 SharedStoreTileTraits;
186 };
187
189
190 template <typename GemmConfig_, typename Index_>
191 struct IgemmTileTraitsHelperA<MatrixLayout::kRowMajor, GemmConfig_, Index_> {
193 static MatrixLayout::Kind const kLayout = MatrixLayout::kRowMajor;
194
196 typedef int8_t Scalar;
198 typedef int8_t MultiplyAddScalar;
199
201 static int const kScalarsPerStsA = 16;
202
204 typedef IgemmGlobalTileTraits<
205 GemmOperand::kA,
206 // The layout.
207 MatrixLayout::kRowMajor,
208 // The pointer is float const.
209 int8_t const,
210 // The tile has size NxK in GEMM's terminology.
211 Shape<1, GemmConfig_::OutputTile::kW, GemmConfig_::OutputTile::kD>,
212 // The threads are distributed as warps x 32 (the traits may reorganize).
213 Shape<1, ShapeCount<typename GemmConfig_::Warps>::kCount, GemmConfig_::kWarpSize>,
214 // The number of scalars per LDG (LDG.32 or LDG.128, etc).
215 GemmConfig_::kScalarsPerLdgA>
216 GlobalTileTraits;
217
219 typedef IgemmGlobalIteratorAb<GlobalTileTraits, Index_> GlobalLoadIterator;
220
222 typedef GemmSharedStoreWithSkewTileAbTraits<
223 // The pointer is int8.
224 int8_t,
225 // The tile has size KxN in GEMM's terminology.
226 Shape<GemmConfig_::kStages, GemmConfig_::OutputTile::kD / 4, GemmConfig_::OutputTile::kW * 4>,
227 // The threads are distributed as (threads / K) x K (the traits may reorganize).
228 typename GlobalTileTraits::Threads,
229 // The number of scalars per STS.
230 kScalarsPerStsA,
231 // The skew to avoid bank conflicts added in the tile W dimension.
232 16>
233 SharedStoreTileTraits;
234
236 typedef GemmSharedLoadTileATraits<
237 // The pointer is float const.
238 int8_t const,
239 // The output tile size.
240 typename GemmConfig_::OutputTile,
241 // The number of warps.
242 typename GemmConfig_::Warps,
243 // The number of threads per warp.
244 typename GemmConfig_::MultiplyAdd::ThreadsPerWarp,
245 // The shape of the FMA instruction.
246 typename GemmConfig_::InstructionShape,
247 // The number of stages.
248 GemmConfig_::kStages,
249 // The number of scalars per LDS.
250 16,
251 // The skew.
252 SharedStoreTileTraits::kSkew>
253 SharedLoadTileTraits;
254 };
255
257
258 template <enum MatrixLayout::Kind kLayout_, typename GemmConfig_, typename Index_>
259 struct IgemmTileTraitsHelperB : public GemmTileTraitsHelperB<kLayout_, GemmConfig_> {};
260
262
263 template <typename GemmConfig_, typename Index_>
264 struct IgemmTileTraitsHelperB<MatrixLayout::kColumnMajor, GemmConfig_, Index_> {
266 static MatrixLayout::Kind const kLayout = MatrixLayout::kColumnMajor;
267
269 typedef int8_t Scalar;
271 typedef int8_t MultiplyAddScalar;
272
274 static int const kScalarsPerStsB = 16;
275
277 typedef IgemmGlobalTileTraits<
278 GemmOperand::kB,
279 // The layout.
280 MatrixLayout::kColumnMajor,
281 // The pointer is float const.
282 int8_t const,
283 // The tile has size NxK in GEMM's terminology.
284 Shape<1, GemmConfig_::OutputTile::kH, GemmConfig_::OutputTile::kD>,
285 // The threads are distributed as warps x 32 (the traits may reorganize).
286 Shape<1, ShapeCount<typename GemmConfig_::Warps>::kCount, GemmConfig_::kWarpSize>,
287 // The number of scalars per LDG (LDG.32 or LDG.128, etc).
288 GemmConfig_::kScalarsPerLdgB>
289 GlobalTileTraits;
290
292 typedef IgemmGlobalIteratorAb<GlobalTileTraits, Index_> GlobalLoadIterator;
293
295 typedef GemmSharedStoreWithSkewTileAbTraits<
296 // The pointer is int8.
297 int8_t,
298 // The tile has size KxN in GEMM's terminology.
299 Shape<GemmConfig_::kStages, GemmConfig_::OutputTile::kD / 4, GemmConfig_::OutputTile::kH * 4>,
300 // The threads are distributed as (threads / K) x K (the traits may reorganize).
301 typename GlobalTileTraits::Threads,
302 // The number of scalars per STS.
303 kScalarsPerStsB,
304 // The skew to avoid bank conflicts added in the tile W dimension.
305 16>
306 SharedStoreTileTraits;
307
309 typedef GemmSharedLoadTileBTraits<
310 // The pointer is float const.
311 int8_t const,
312 // The output tile size.
313 typename GemmConfig_::OutputTile,
314 // The number of warps.
315 typename GemmConfig_::Warps,
316 // The number of threads per warp.
317 typename GemmConfig_::MultiplyAdd::ThreadsPerWarp,
318 // The shape of the FMA instruction.
319 typename GemmConfig_::InstructionShape,
320 // The number of stages.
321 GemmConfig_::kStages,
322 // The number of scalars per LDS.
323 16,
324 // The skew.
325 SharedStoreTileTraits::kSkew>
326 SharedLoadTileTraits;
327 };
328
330
331 template <typename GemmConfig_, typename Index_>
332 struct IgemmTileTraitsHelperB<MatrixLayout::kRowMajor, GemmConfig_, Index_>
333 : public GemmTileTraitsHelperB<MatrixLayout::kRowMajor, GemmConfig_> {
335 typedef GemmTileTraitsHelperB<MatrixLayout::kRowMajor, GemmConfig_> Base;
336
338 static int const kScalarsPerStsB = 16;
339
341 typedef IgemmGlobalTileTraits<
342 GemmOperand::kB,
343 // The layout.
344 MatrixLayout::kRowMajor,
345 // The pointer is float const.
346 int8_t const,
347 // The tile has size KxM in GEMM's terminology.
348 Shape<1, GemmConfig_::OutputTile::kD, GemmConfig_::OutputTile::kH>,
349 // The threads are distributed as warps x 32 (the traits may reorganize).
350 Shape<1, ShapeCount<typename GemmConfig_::Warps>::kCount, GemmConfig_::kWarpSize>,
351 // The number of scalars per LDG (LDG.32 or LDG.128, etc).
352 GemmConfig_::kScalarsPerLdgB>
353 GlobalTileTraits;
354
356 typedef GemmGlobalIteratorAb<GlobalTileTraits, Index_> GlobalLoadIterator;
357
359 typedef GemmSharedStoreTileAbTraits<
360 // The pointer is float.
361 int8_t,
362 // The tile has size KxM in GEMM's terminology.
363 Shape<GemmConfig_::kStages, GemmConfig_::OutputTile::kD / 4, GemmConfig_::OutputTile::kH * 4>,
364 // The threads are distributed as warps x 32 (the traits may reorganize).
365 typename GlobalTileTraits::Threads,
366 // The number of scalars per STS (STS.32 or STS.128, etc).
367 kScalarsPerStsB>
368 SharedStoreTileTraits;
369 };
370
372
373 template <enum MatrixLayout::Kind kLayout_, typename Iterator_>
374 struct IgemmTransformerA {};
375
376 template <typename Iterator_>
377 struct IgemmTransformerA<MatrixLayout::kRowMajor, Iterator_> {
378 typedef Copy<typename Iterator_::Fragment> Transformer;
379 };
380
381 template <typename Iterator_>
382 struct IgemmTransformerA<MatrixLayout::kColumnMajor, Iterator_> {
383 typedef IgemmSwizzle<Iterator_> Transformer;
384 };
385
387
388 template <enum MatrixLayout::Kind kLayout_, typename Iterator_>
389 struct IgemmTransformerB {};
390
391 template <typename Iterator_>
392 struct IgemmTransformerB<MatrixLayout::kColumnMajor, Iterator_> {
393 typedef Copy<typename Iterator_::Fragment> Transformer;
394 };
395
396 template <typename Iterator_>
397 struct IgemmTransformerB<MatrixLayout::kRowMajor, Iterator_> {
398 typedef IgemmSwizzle<Iterator_> Transformer;
399 };
400
402
403 template <
405 MatrixLayout::Kind kLayoutA_,
407 MatrixLayout::Kind kLayoutB_,
409 typename OutputTile_,
411 typename ScalarD_,
413 typename EpilogueFunctor_,
415 typename ThreadGemmShape_ = Shape<32, 8, 8>,
417 typename Index_ = int>
418 struct IgemmTraitsHelper {
420 typedef IgemmConfig<OutputTile_, ScalarD_, ThreadGemmShape_> GemmConfig;
422 typedef IgemmTileTraitsHelperA<kLayoutA_, GemmConfig, Index_> GemmTileTraitsHelperA;
424 typedef IgemmTileTraitsHelperB<kLayoutB_, GemmConfig, Index_> GemmTileTraitsHelperB;
425
427 typedef typename GemmTileTraitsHelperA::GlobalLoadIterator GlobalLoadIteratorA;
429 typedef typename IgemmTransformerA<GemmTileTraitsHelperA::kLayout,
430 GlobalLoadIteratorA>::Transformer GlobalTransformerA;
432 typedef TileStoreIterator<typename GemmTileTraitsHelperA::SharedStoreTileTraits,
433 typename GemmTileTraitsHelperA::SharedStoreTileTraits::Scalar,
434 IteratorAdvance::kH,
435 MemorySpace::kShared>
436 SharedStoreIteratorA;
438 typedef GlobalLoadStream<GemmOperand::kA,
439 GlobalLoadIteratorA,
440 SharedStoreIteratorA,
441 GlobalTransformerA>
442 GlobalLoadStreamA;
443
445 typedef typename GemmTileTraitsHelperB::GlobalLoadIterator GlobalLoadIteratorB;
446 // The default transformer for B.
447 typedef typename IgemmTransformerB<GemmTileTraitsHelperB::kLayout,
448 GlobalLoadIteratorB>::Transformer GlobalTransformerB;
450 typedef TileStoreIterator<typename GemmTileTraitsHelperB::SharedStoreTileTraits,
451 typename GemmTileTraitsHelperB::SharedStoreTileTraits::Scalar,
452 IteratorAdvance::kH,
453 MemorySpace::kShared>
454 SharedStoreIteratorB;
456 typedef GlobalLoadStream<GemmOperand::kB,
457 GlobalLoadIteratorB,
458 SharedStoreIteratorB,
459 GlobalTransformerB>
460 GlobalLoadStreamB;
461
463 typedef TileLoadIterator<typename GemmTileTraitsHelperA::SharedLoadTileTraits,
464 typename GemmTileTraitsHelperA::SharedLoadTileTraits::Scalar,
465 IteratorAdvance::kH,
466 MemorySpace::kShared>
467 SharedLoadIteratorA;
469 typedef SharedLoadStream<SharedLoadIteratorA, Copy<typename SharedLoadIteratorA::Fragment> >
470 SharedLoadStreamA;
472 typedef TileLoadIterator<typename GemmTileTraitsHelperB::SharedLoadTileTraits,
473 typename GemmTileTraitsHelperB::SharedLoadTileTraits::Scalar,
474 IteratorAdvance::kH,
475 MemorySpace::kShared>
476 SharedLoadIteratorB;
478 typedef SharedLoadStream<SharedLoadIteratorB, Copy<typename SharedLoadIteratorB::Fragment> >
479 SharedLoadStreamB;
480
482 typedef typename GemmConfig::MultiplyAdd MultiplyAdd;
484 typedef ClearAccumulators<typename MultiplyAdd::ScalarC> ClearAccumulators;
485
487 typedef IgemmEpilogue<IgemmEpilogueTraits<GemmConfig, EpilogueFunctor_> > Epilogue;
488 };
489
491
492 template <typename ScalarD_>
493 struct IgemmEpilogueScalar {
494 typedef float Scalar;
495 };
496
497 template <>
498 struct IgemmEpilogueScalar<int> {
499 typedef int Scalar;
500 };
501
503
504 template <
506 MatrixLayout::Kind kLayoutA_,
508 MatrixLayout::Kind kLayoutB_,
510 typename OutputTile_ = Shape<32, 128, 128>,
512 typename ScalarD_ = int,
514 typename EpilogueFunctor_ = LinearScaling<typename IgemmEpilogueScalar<ScalarD_>::Scalar>,
516 typename ThreadGemmShape_ = Shape<32, 8, 8>,
518 typename Index_ = int,
520 typename Helper_ = IgemmTraitsHelper<kLayoutA_,
521 kLayoutB_,
522 OutputTile_,
523 ScalarD_,
524 EpilogueFunctor_,
525 ThreadGemmShape_,
526 Index_> >
527 struct IgemmTraits : public GemmTraits<
528 // The config.
529 typename Helper_::GemmConfig,
530 // The stream to load A from global memory to shared memory.
531 typename Helper_::GlobalLoadStreamA,
532 // The stream to load B from global memory to shared memory.
533 typename Helper_::GlobalLoadStreamB,
534 // The stream to load A from shared memory.
535 typename Helper_::SharedLoadStreamA,
536 // The stream to load B from shared memory.
537 typename Helper_::SharedLoadStreamB,
538 // The epilogue.
539 typename Helper_::Epilogue,
540 // The block swizzle to reorganize the grid.
541 IdentityBlockSwizzle,
542 // The index.
543 Index_,
544 // The tool used to clear accumulators.
545 typename Helper_::ClearAccumulators> {};
546
548
549 } // namespace gemm
550 } // namespace cutlass
cutlass::gemm::IgemmTraitsHelper::GlobalTransformerB
IgemmTransformerB< GemmTileTraitsHelperB::kLayout, GlobalLoadIteratorB >::Transformer GlobalTransformerB
Definition: igemm_traits.h:448
cutlass::MemorySpace::kShared
Definition: load_store.h:41
cutlass::gemm::IgemmTileTraitsHelperB< MatrixLayout::kRowMajor, GemmConfig_, Index_ >::Base
GemmTileTraitsHelperB< MatrixLayout::kRowMajor, GemmConfig_ > Base
The base config.
Definition: igemm_traits.h:335
cutlass
Definition: convert.h:33
cutlass::gemm::GemmSharedLoadTileATraits
Definition: gemm_shared_tile.h:128
cutlass::gemm::IgemmGlobalTileTraits::Threads
Base::Threads Threads
The threads.
Definition: igemm_global_tile.h:66
cutlass::gemm::IgemmTraitsHelper::GemmTileTraitsHelperB
IgemmTileTraitsHelperB< kLayoutB_, GemmConfig, Index_ > GemmTileTraitsHelperB
The GEMM config for B.
Definition: igemm_traits.h:424
cutlass::gemm::GemmSharedStoreWithSkewTileAbTraits
Definition: gemm_shared_tile.h:80
cutlass::gemm::IgemmTransformerA< MatrixLayout::kColumnMajor, Iterator_ >::Transformer
IgemmSwizzle< Iterator_ > Transformer
Definition: igemm_traits.h:383
gemm_global_tile.h
Defines iterators for efficiently loading and storing to global memory.
igemm_swizzle.h
Transposes a fragment of data containing packed 8-bit integer elements.
cutlass::gemm::IgemmTransformerB< MatrixLayout::kColumnMajor, Iterator_ >::Transformer
Copy< typename Iterator_::Fragment > Transformer
Definition: igemm_traits.h:393
cutlass::gemm::IgemmTileTraitsHelperA< MatrixLayout::kRowMajor, GemmConfig_, Index_ >::SharedStoreTileTraits
GemmSharedStoreWithSkewTileAbTraits< int8_t, Shape< GemmConfig_::kStages, GemmConfig_::OutputTile::kD/4, GemmConfig_::OutputTile::kW *4 >, typename GlobalTileTraits::Threads, kScalarsPerStsA, 16 > SharedStoreTileTraits
The traits class to build the iterator to store data to shared memory for A^N.
Definition: igemm_traits.h:233
cutlass::gemm::IgemmTileTraitsHelperB< MatrixLayout::kColumnMajor, GemmConfig_, Index_ >::GlobalTileTraits
IgemmGlobalTileTraits< GemmOperand::kB, MatrixLayout::kColumnMajor, int8_t const, Shape< 1, GemmConfig_::OutputTile::kH, GemmConfig_::OutputTile::kD >, Shape< 1, ShapeCount< typename GemmConfig_::Warps >::kCount, GemmConfig_::kWarpSize >, GemmConfig_::kScalarsPerLdgB > GlobalTileTraits
The traits class to build the iterator to load data from global memory for B^T.
Definition: igemm_traits.h:289
gemm_traits.h
Defines structural properties of complete GEMM computation.
cutlass::gemm::IgemmTileTraitsHelperA< MatrixLayout::kRowMajor, GemmConfig_, Index_ >::GlobalLoadIterator
IgemmGlobalIteratorAb< GlobalTileTraits, Index_ > GlobalLoadIterator
The global load iterator.
Definition: igemm_traits.h:219
cutlass::gemm::IgemmTileTraitsHelperA
Definition: igemm_traits.h:144
cutlass::gemm::IgemmEpilogue
Definition: igemm_epilogue.h:290
cutlass::MatrixLayout::Kind
Kind
Enumeration defining fundamental contiguous layouts.
Definition: matrix_traits.h:159
cutlass::Copy
Definition: convert.h:69
cutlass::gemm::IgemmTileTraitsHelperB< MatrixLayout::kRowMajor, GemmConfig_, Index_ >::GlobalTileTraits
IgemmGlobalTileTraits< GemmOperand::kB, MatrixLayout::kRowMajor, int8_t const, Shape< 1, GemmConfig_::OutputTile::kD, GemmConfig_::OutputTile::kH >, Shape< 1, ShapeCount< typename GemmConfig_::Warps >::kCount, GemmConfig_::kWarpSize >, GemmConfig_::kScalarsPerLdgB > GlobalTileTraits
The traits class to build the iterator to load data from global memory for B^T.
Definition: igemm_traits.h:353
cutlass::gemm::GemmSharedStoreTileAbTraits
Definition: gemm_shared_tile.h:38
cutlass::IteratorAdvance::kH
Definition: tile_iterator.h:65
cutlass::gemm::IgemmTileTraitsHelperA< MatrixLayout::kRowMajor, GemmConfig_, Index_ >::MultiplyAddScalar
int8_t MultiplyAddScalar
The scalar stored in shared memory.
Definition: igemm_traits.h:198
cutlass::gemm::IgemmTraitsHelper::GlobalLoadIteratorB
GemmTileTraitsHelperB::GlobalLoadIterator GlobalLoadIteratorB
The iterator to load B from global memory.
Definition: igemm_traits.h:445
igemm_multiply_add.h
Implements matrix multiply accumulate operation of 8-bit integer data using DP4A instruction.
cutlass::gemm::GemmSharedLoadTileBTraits
Definition: gemm_shared_tile.h:200
cutlass::gemm::IgemmTraitsHelper::SharedStoreIteratorB
TileStoreIterator< typename GemmTileTraitsHelperB::SharedStoreTileTraits, typename GemmTileTraitsHelperB::SharedStoreTileTraits::Scalar, IteratorAdvance::kH, MemorySpace::kShared > SharedStoreIteratorB
The iterator to store B to shared memory.
Definition: igemm_traits.h:454
cutlass::gemm::IgemmTileTraitsHelperB< MatrixLayout::kColumnMajor, GemmConfig_, Index_ >::SharedLoadTileTraits
GemmSharedLoadTileBTraits< int8_t const, typename GemmConfig_::OutputTile, typename GemmConfig_::Warps, typename GemmConfig_::MultiplyAdd::ThreadsPerWarp, typename GemmConfig_::InstructionShape, GemmConfig_::kStages, 16, SharedStoreTileTraits::kSkew > SharedLoadTileTraits
The traits class to build the iterator to load from shared memory for B^N.
Definition: igemm_traits.h:326
cutlass::gemm::GemmGlobalIteratorAb
Definition: gemm_global_tile.h:163
cutlass::gemm::IgemmTileTraitsHelperB< MatrixLayout::kColumnMajor, GemmConfig_, Index_ >::MultiplyAddScalar
int8_t MultiplyAddScalar
The scalar stored in shared memory.
Definition: igemm_traits.h:271
gemm_epilogue.h
Implements the epilogue phase of the GEMM kernel that efficiently updates global memory with the comp...
cutlass::gemm::IgemmTileTraitsHelperA< MatrixLayout::kRowMajor, GemmConfig_, Index_ >::GlobalTileTraits
IgemmGlobalTileTraits< GemmOperand::kA, MatrixLayout::kRowMajor, int8_t const, Shape< 1, GemmConfig_::OutputTile::kW, GemmConfig_::OutputTile::kD >, Shape< 1, ShapeCount< typename GemmConfig_::Warps >::kCount, GemmConfig_::kWarpSize >, GemmConfig_::kScalarsPerLdgA > GlobalTileTraits
The traits class to build the iterator to load data from global memory for A^T.
Definition: igemm_traits.h:216
cutlass::gemm::GlobalLoadStream
Definition: gemm_global_stream.h:52
cutlass::gemm::GemmTileTraitsHelperB
Definition: gemm_traits.h:191
cutlass::gemm::IgemmTraitsHelper::Epilogue
IgemmEpilogue< IgemmEpilogueTraits< GemmConfig, EpilogueFunctor_ > > Epilogue
The epilogue.
Definition: igemm_traits.h:487
cutlass::gemm::IgemmEpilogueScalar< int >::Scalar
int Scalar
Definition: igemm_traits.h:499
cutlass::gemm::IgemmTransformerB< MatrixLayout::kRowMajor, Iterator_ >::Transformer
IgemmSwizzle< Iterator_ > Transformer
Definition: igemm_traits.h:398
cutlass::gemm::IgemmTileTraitsHelperA< MatrixLayout::kColumnMajor, GemmConfig_, Index_ >::SharedStoreTileTraits
GemmSharedStoreTileAbTraits< int8_t, Shape< GemmConfig_::kStages, GemmConfig_::OutputTile::kD/4, GemmConfig_::OutputTile::kW *4 >, typename GlobalTileTraits::Threads, kScalarsPerStsA > SharedStoreTileTraits
The traits class to build the iterator to store data to shared memory for A^N.
Definition: igemm_traits.h:185
cutlass::gemm::IgemmSwizzle
Definition: igemm_swizzle.h:38
cutlass::gemm::IgemmTileTraitsHelperB
Definition: igemm_traits.h:259
cutlass::gemm::IgemmTraitsHelper
Definition: igemm_traits.h:418
cutlass::TileLoadIterator
An iterator implementing Tile Load Iterator Concept for loading a tile from memory.
Definition: tile_iterator.h:402
cutlass::gemm::IgemmTraitsHelper::GlobalTransformerA
IgemmTransformerA< GemmTileTraitsHelperA::kLayout, GlobalLoadIteratorA >::Transformer GlobalTransformerA
The default transformer for A.
Definition: igemm_traits.h:430
gemm_shared_tile.h
Defines iterators for efficiently loading and storing tiles to and from shared memory.
cutlass::gemm::IgemmTraitsHelper::GlobalLoadStreamB
GlobalLoadStream< GemmOperand::kB, GlobalLoadIteratorB, SharedStoreIteratorB, GlobalTransformerB > GlobalLoadStreamB
The stream to load B from global memory to shared memory.
Definition: igemm_traits.h:460
cutlass::gemm::SharedLoadStream
Definition: gemm_shared_stream.h:45
cutlass::gemm::IgemmGlobalTileTraits
Definition: igemm_global_tile.h:50
reshape_tile.h
Defines a type for restructuring a tile.
cutlass::gemm::IgemmTraitsHelper::GlobalLoadIteratorA
GemmTileTraitsHelperA::GlobalLoadIterator GlobalLoadIteratorA
The iterator to load A from global memory.
Definition: igemm_traits.h:427
cutlass::gemm::GemmConfig
Definition: gemm_config.h:76
cutlass::gemm::IgemmTraitsHelper::SharedStoreIteratorA
TileStoreIterator< typename GemmTileTraitsHelperA::SharedStoreTileTraits, typename GemmTileTraitsHelperA::SharedStoreTileTraits::Scalar, IteratorAdvance::kH, MemorySpace::kShared > SharedStoreIteratorA
The iterator to store A to shared memory.
Definition: igemm_traits.h:436
cutlass::gemm::GemmTileTraitsHelperA
Definition: gemm_traits.h:52
cutlass::GemmOperand::kB
Definition: matrix_traits.h:357
cutlass::gemm::IgemmConfig
Definition: igemm_traits.h:57
cutlass::gemm::IgemmTileTraitsHelperA< MatrixLayout::kColumnMajor, GemmConfig_, Index_ >::GlobalTileTraits
IgemmGlobalTileTraits< GemmOperand::kA, MatrixLayout::kColumnMajor, int8_t const, Shape< 1, GemmConfig_::OutputTile::kD, GemmConfig_::OutputTile::kW >, Shape< 1, ShapeCount< typename GemmConfig_::Warps >::kCount, GemmConfig_::kWarpSize >, GemmConfig_::kScalarsPerLdgA > GlobalTileTraits
The traits class to build the iterator to load data from global memory for A^N.
Definition: igemm_traits.h:170
cutlass::gemm::IgemmGlobalIteratorAb
Definition: igemm_global_tile.h:95
cutlass::gemm::IgemmTransformerA
Definition: igemm_traits.h:374
cutlass::gemm::IgemmEpilogueScalar::Scalar
float Scalar
Definition: igemm_traits.h:494
cutlass::gemm::GemmTraits
Definition: gemm_traits.h:349
cutlass::gemm::IgemmTransformerA< MatrixLayout::kRowMajor, Iterator_ >::Transformer
Copy< typename Iterator_::Fragment > Transformer
Definition: igemm_traits.h:378
cutlass::gemm::IgemmTraits
Definition: igemm_traits.h:527
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::IgemmTraitsHelper::SharedLoadIteratorB
TileLoadIterator< typename GemmTileTraitsHelperB::SharedLoadTileTraits, typename GemmTileTraitsHelperB::SharedLoadTileTraits::Scalar, IteratorAdvance::kH, MemorySpace::kShared > SharedLoadIteratorB
The iterator to load B from shared memory.
Definition: igemm_traits.h:476
cutlass::gemm::IgemmTileTraitsHelperB< MatrixLayout::kRowMajor, GemmConfig_, Index_ >::SharedStoreTileTraits
GemmSharedStoreTileAbTraits< int8_t, Shape< GemmConfig_::kStages, GemmConfig_::OutputTile::kD/4, GemmConfig_::OutputTile::kH *4 >, typename GlobalTileTraits::Threads, kScalarsPerStsB > SharedStoreTileTraits
The traits class to build the iterator to store data to shared memory for B^N.
Definition: igemm_traits.h:368
cutlass::MatrixLayout::kRowMajor
Definition: matrix_traits.h:159
cutlass::gemm::GemmGlobalTileTraits::Threads
ReshapeThreads< VectorizedTile, Threads_ >::Threads Threads
The threads shape.
Definition: gemm_global_tile.h:88
cutlass::gemm::GemmTileTraitsHelperB< MatrixLayout::kRowMajor, GemmConfig_ >
Definition: gemm_traits.h:267
cutlass::gemm::ThreadMultiplyAdd
Template performing matrix multiply-add operation within a thread.
Definition: thread_multiply_add.h:44
cutlass::gemm::GemmTileTraitsHelperA< MatrixLayout::kColumnMajor, GemmConfig_ >
Definition: gemm_traits.h:57
cutlass::gemm::IgemmTileTraitsHelperA< MatrixLayout::kRowMajor, GemmConfig_, Index_ >::Scalar
int8_t Scalar
The input scalar.
Definition: igemm_traits.h:196
cutlass::gemm::IgemmTraitsHelper::GemmConfig
IgemmConfig< OutputTile_, ScalarD_, ThreadGemmShape_ > GemmConfig
The IGEMM config.
Definition: igemm_traits.h:420
cutlass::gemm::IgemmTileTraitsHelperB< MatrixLayout::kColumnMajor, GemmConfig_, Index_ >::GlobalLoadIterator
IgemmGlobalIteratorAb< GlobalTileTraits, Index_ > GlobalLoadIterator
The global load iterator.
Definition: igemm_traits.h:292
cutlass::gemm::IgemmTileTraitsHelperA< MatrixLayout::kColumnMajor, GemmConfig_, Index_ >::GlobalLoadIterator
GemmGlobalIteratorAb< GlobalTileTraits, Index_ > GlobalLoadIterator
The global load iterator.
Definition: igemm_traits.h:173
cutlass::gemm::IgemmTileTraitsHelperB< MatrixLayout::kRowMajor, GemmConfig_, Index_ >::GlobalLoadIterator
GemmGlobalIteratorAb< GlobalTileTraits, Index_ > GlobalLoadIterator
The global load iterator.
Definition: igemm_traits.h:356
cutlass::gemm::IgemmTraitsHelper::MultiplyAdd
GemmConfig::MultiplyAdd MultiplyAdd
The multiply-add functor.
Definition: igemm_traits.h:482
cutlass::gemm::IgemmTransformerB
Definition: igemm_traits.h:389
cutlass::gemm::LinearScaling
Functor to compute linear combination of fragments.
Definition: linear_scaling.h:51
cutlass::gemm::IgemmTraitsHelper::SharedLoadStreamA
SharedLoadStream< SharedLoadIteratorA, Copy< typename SharedLoadIteratorA::Fragment > > SharedLoadStreamA
The stream to load A from shared memory.
Definition: igemm_traits.h:470
cutlass::GemmOperand::kA
Definition: matrix_traits.h:357
cutlass::gemm::IgemmTraitsHelper::GlobalLoadStreamA
GlobalLoadStream< GemmOperand::kA, GlobalLoadIteratorA, SharedStoreIteratorA, GlobalTransformerA > GlobalLoadStreamA
The stream to load A from global memory to shared memory.
Definition: igemm_traits.h:442
cutlass::gemm::IgemmTraitsHelper::GemmTileTraitsHelperA
IgemmTileTraitsHelperA< kLayoutA_, GemmConfig, Index_ > GemmTileTraitsHelperA
The GEMM config for A.
Definition: igemm_traits.h:422
gemm.h
Implements a software-pipelined efficient GEMM.
cutlass::gemm::IgemmTileTraitsHelperA< MatrixLayout::kRowMajor, GemmConfig_, Index_ >::SharedLoadTileTraits
GemmSharedLoadTileATraits< int8_t const, typename GemmConfig_::OutputTile, typename GemmConfig_::Warps, typename GemmConfig_::MultiplyAdd::ThreadsPerWarp, typename GemmConfig_::InstructionShape, GemmConfig_::kStages, 16, SharedStoreTileTraits::kSkew > SharedLoadTileTraits
The traits class to build the iterator to load from shared memory for A^N.
Definition: igemm_traits.h:253
cutlass::gemm::IgemmTraitsHelper::SharedLoadStreamB
SharedLoadStream< SharedLoadIteratorB, Copy< typename SharedLoadIteratorB::Fragment > > SharedLoadStreamB
The stream to load B from shared memory.
Definition: igemm_traits.h:479
gemm_epilogue_traits.h
Defines structural properties of the GEMM epilogue.
cutlass::gemm::IgemmEpilogueScalar
Definition: igemm_traits.h:493
igemm_epilogue.h
Defines the epilogue phase of the GEMM computation for IGEMM, supporting integer and floating-point o...
convert.h
Defines conversion operations among Fragments of different base type.
cutlass::gemm::IgemmTileTraitsHelperB< MatrixLayout::kColumnMajor, GemmConfig_, Index_ >::SharedStoreTileTraits
GemmSharedStoreWithSkewTileAbTraits< int8_t, Shape< GemmConfig_::kStages, GemmConfig_::OutputTile::kD/4, GemmConfig_::OutputTile::kH *4 >, typename GlobalTileTraits::Threads, kScalarsPerStsB, 16 > SharedStoreTileTraits
The traits class to build the iterator to store data to shared memory for B^N.
Definition: igemm_traits.h:306
cutlass::gemm::IgemmTileTraitsHelperB< MatrixLayout::kColumnMajor, GemmConfig_, Index_ >::Scalar
int8_t Scalar
The input scalar.
Definition: igemm_traits.h:269
igemm_global_tile.h
Implements tile iterators to partition the thread block tile into 2D subtiles and efficiently load ea...
cutlass::gemm::IgemmTraitsHelper::SharedLoadIteratorA
TileLoadIterator< typename GemmTileTraitsHelperA::SharedLoadTileTraits, typename GemmTileTraitsHelperA::SharedLoadTileTraits::Scalar, IteratorAdvance::kH, MemorySpace::kShared > SharedLoadIteratorA
The iterator to load A from shared memory.
Definition: igemm_traits.h:467
cutlass::gemm::IgemmTileTraitsHelperA< MatrixLayout::kColumnMajor, GemmConfig_, Index_ >::Base
GemmTileTraitsHelperA< MatrixLayout::kColumnMajor, GemmConfig_ > Base
The base config.
Definition: igemm_traits.h:152
cutlass::TileStoreIterator
An iterator implementing Tile Store Iterator Concept for storing a tile to memory.
Definition: tile_iterator.h:841
cutlass::gemm::IgemmTraitsHelper::ClearAccumulators
ClearAccumulators< typename MultiplyAdd::ScalarC > ClearAccumulators
The object to clear accumulators.
Definition: igemm_traits.h:484