Compare commits

...

340 Commits

Author SHA1 Message Date
c4f6b8c6bc Updates for 3.0 (#857)
Co-authored-by: Aniket Shivam <ashivam@nvidia.com>
2023-03-09 15:27:40 -05:00
a68e2f95f0 Reduce versbosity in manifest.py (#845) 2023-03-07 11:53:01 -05:00
a31b43b3f3 Re-enable aarch64 support lost in 277bd6e537 (#846) 2023-03-02 11:17:21 -05:00
f396cdd15c ex24[gemm_grouped]: Allow to change layout/dtype (#841)
* ex24[gemm_grouped]: Allow to change layout/dtype

* Address suggestion from @jackkosaian

---------

Co-authored-by: danthe3rd <danthe3rd>
2023-03-01 07:13:51 -05:00
92ebbf1dc4 Fix typos (#839) 2023-02-27 11:17:58 -05:00
65688c2a87 streamk fix (#836)
Co-authored-by: Haicheng Wu <haichengw@nvidia.com>
2023-02-23 16:35:08 -05:00
f303889ed9 fMHA: Sync FW with xFormers (#828)
* fMHA: Add support for bias+dropout in FW

* Remove 'getMaximumSharedMemoryPerBlockKb'

* fix comments

---------

Co-authored-by: danthe3rd <danthe3rd>
Co-authored-by: Haicheng Wu <haichengw@nvidia.com>
2023-02-22 23:25:31 -05:00
9cdbe33570 Add fixed_channel and few_channel mode to int8 in generator (#829) 2023-02-21 21:15:39 -05:00
95f673ecf7 Update base_grouped.h (#832) 2023-02-21 14:48:30 -05:00
91b8de8d32 streamk fix (#830)
Co-authored-by: Haicheng Wu <haichengw@nvidia.com>
2023-02-20 11:03:16 -05:00
d8359c804b Changes to iterators to support s8 gemm with f16 outputs (#812)
* Changes to iterators to support s8 gemm with f16 outputs

* should work

---------

Co-authored-by: Sujan Gonugondla <gsujan@amaon.com>
Co-authored-by: Haicheng Wu <haichengw@nvidia.com>
2023-02-16 18:37:51 -05:00
34bed24af3 Update helper.h
copyright banner
2023-02-16 16:50:04 -05:00
ZZK
a101ac283f Fix some typos (#791)
* fix typo

* fix a deadlink to code
2023-02-16 15:56:55 -05:00
9fb38ac048 fix alignmentC=8 for imma N=128 (#822)
Co-authored-by: Haicheng Wu <haichengw@nvidia.com>
2023-02-15 12:06:00 -05:00
8f5c242426 Update dual_gemm_common.h
fix the copyright of a new file.
2023-02-13 15:35:33 -05:00
3c995c7606 Extend DualGemm: support batched mode + decouple B0/B1 layouts (#790)
* Fix MHA kernel

Summary:

ATT

Test Plan:

Reviewers:

Subscribers:

Tasks:

Tags:

* Extend DualGemm to support batched mode (#5)

Following the GemmUniversalMode::kBatched implementation, batched mode is added to the DualGemm (under examples/45_dual_gemm). DualGemmMode::kBatched and SplitKSerial are not compatible: Status::kErrorInvalidProblem is returned if both are set.

* Decouple LayoutB0 and LayoutB1 in DualGemm

The DualGemm template assumed the same layout, LayoutB, for both right operand matrices B0 and B1. This is problematic if the layout of the two matrices is different. In particular, this may be the case when one of the matrices is row-major, while the other is a (column) vector that has to be broadcasted in column-major with zero stride (e.g., as {B1.device_data(), 0}) for the DualGemm implementation to be able to process B0 and B1 simultaneously.

In this commit, LayoutB0 and LayoutB1 are decoupled throughout the DualGemm code (device, kernel, and mma). Additionally, the batch strides of B0 and B1 are also decoupled to accommodate the column vector B1 case described above.

* Remove comment as no longer relevant

* Revert Fix MHA kernel

---------

Co-authored-by: mikeiovine <mikeiovine@fb.com>
2023-02-13 15:27:13 -05:00
ce8597dc14 Fix type bug in conv2d/gemm with broadcast (#796)
add ElementVector

---------

Co-authored-by: Haicheng Wu <haichengw@nvidia.com>
2023-02-09 20:53:25 -05:00
2e10404d26 xFormer updates to fMHA FW (#773)
* xFormer updates to fMHA FW

* Convert format to BMHK for '41_fused_multi_head_attention_fixed_seqlen'

* Add missing files

* Remove xFormers specific code

* Update fused_multihead_attention_fixed_seqlen.cu

* rebase and solve conflicts

* remove white space

---------

Co-authored-by: danthe3rd <danthe3rd>
Co-authored-by: Haicheng Wu <haichengw@nvidia.com>
2023-02-08 23:00:10 -05:00
5ff5209ed5 Add acc2smem in epilogue/threadblock/epilogue.h (#806) 2023-02-06 22:04:16 -05:00
5921043981 Re-enable all alignments for int accumulators (#807) 2023-02-06 22:01:15 -05:00
add4ba622f Fix 8.4 + CUDA 11.4 build (#789)
Work around a likely GCC 8.x issue with fold expressions
and generic lambdas.

Only use the work-around when the host compiler is GCC 8.x.
This avoids any concerns about the work-around possibly
hindering inlining for a critical CuTe function (product).

Users can experiment with the work-around for other compilers
or compiler versions by defining the following macro.

CUTE_FOLD_GENERIC_LAMBDA_WORKAROUND

Fixes https://github.com/NVIDIA/cutlass/issues/788

Co-authored-by: Mark Hoemmen <mhoemmen@nvidia.com>
2023-01-27 09:18:59 -05:00
277bd6e537 CUTLASS 3.0.0 (#786)
* CUTLASS 3.0.0
2023-01-23 20:55:28 -05:00
66d9cddc83 New updates for 2.11 (#775)
* New updates.

* Minor profiler updates

Co-authored-by: Aniket Shivam <ashivam@nvidia.com>
2023-01-20 16:32:57 -05:00
d49bef88f9 Enable aarch64 support (#779) 2023-01-20 15:51:58 -05:00
8b42e751c6 streamk paper link (#765)
Co-authored-by: Haicheng Wu <haichengw@nvidia.com>
2023-01-10 22:10:43 -05:00
eb7f99d3dd @hwu36 Adding the individual arXiv link for Stream-K paper. (#764)
* Stream-K individual paper entry.

* arXiv links updated.
2023-01-10 20:39:06 -05:00
764b840d6f streamk example and performance tuning (#760)
* streamk example and performance tuning

* one missing file

Co-authored-by: Haicheng Wu <haichengw@nvidia.com>
2023-01-10 16:10:02 -05:00
a1046d49c1 Adds missing semicolon (#759) 2023-01-09 21:50:46 -05:00
1cd994b4cf Update PUBLICATIONS.md
@neoblizz @dumerrill 

thesis covering streamk
2023-01-08 00:42:19 -05:00
7bdba07310 Add definitions for tag structs. (#752)
This commit changes the declarations of MMA operator class (SIMT, Tensor Core, WMMA Tensor Core) and operator type (multiply-add and so on) to definitions. This is done so that these tag structs are no longer incomplete types, which allows the `typeid` operator to be used on these tag structs. This is necessary for these tag structs to be used as type parameters in [GoogleTest typed tests](https://google.github.io/googletest/advanced.html#typed-tests).
2023-01-06 09:46:52 -05:00
c54ede3a9e Add const overloads for iterator functions. (#753)
This commit adds `const`-correct overloads for `Array::{begin,end,rbegin,rend}`. These overloads are necessary for usage with [the GMock Container Matchers](http://google.github.io/googletest/reference/matchers.html#container-matchers), which cast the `Container` argument to a constant reference.
2023-01-06 09:46:34 -05:00
ff6e733fe1 restore the old epilogue for everything except streamk (#749)
Co-authored-by: Haicheng Wu <haichengw@nvidia.com>
2023-01-04 11:02:55 -05:00
5989b7e1d7 Update PUBLICATIONS.md
Add coconet paper to the publication list.  @abhijangda
2023-01-04 09:18:38 -05:00
1e64f153b3 improve streamk load balance (#743)
Co-authored-by: Haicheng Wu <haichengw@nvidia.com>
2022-12-25 13:56:33 -05:00
78b30d3191 Update README.md 2022-12-21 11:58:19 -05:00
59de82688b Update README.md 2022-12-21 11:57:55 -05:00
b85865d1ad Add missing #include directives (#741)
This commit adds two `#include` directives so that the definitions of `cutlass::gemm::warp::WarpSize` from "cutlass/gemm/warp/mma.h" and `cutlass::arch::OpClassSimt` from "cutlass/arch/mma.h" are visible to "cutlass/epilogue/threadblock/default_epilogue_simt.h". Without them, there are compiler errors when building the header standalone:

```
In file included from cutlass/include/cutlass/epilogue/threadblock/default_epilogue_simt.cu:1:
./cutlass/include/cutlass/epilogue/threadblock/default_epilogue_simt.h:351:32: error: no member named 'warp' in namespace 'cutlass::gemm'; did you mean simply 'warp'?
  static int const kWarpSize = cutlass::gemm::warp::WarpSize<arch::OpClassSimt>::value;
                               ^
./cutlass/include/cutlass/epilogue/warp/tile_iterator_simt.h:49:11: note: 'warp' declared here
namespace warp {
          ^
In file included from cutlass/include/cutlass/epilogue/threadblock/default_epilogue_simt.cu:1:
./cutlass/include/cutlass/epilogue/threadblock/default_epilogue_simt.h:351:53: error: no member named 'WarpSize' in namespace 'cutlass::epilogue::warp'
  static int const kWarpSize = cutlass::gemm::warp::WarpSize<arch::OpClassSimt>::value;
                                              ~~~~~~^
./cutlass/include/cutlass/epilogue/threadblock/default_epilogue_simt.h:351:68: error: no member named 'OpClassSimt' in namespace 'cutlass::arch'
  static int const kWarpSize = cutlass::gemm::warp::WarpSize<arch::OpClassSimt>::value;
                                                             ~~~~~~^
./cutlass/include/cutlass/epilogue/threadblock/default_epilogue_simt.h:351:82: error: no member named 'value' in the global namespace
  static int const kWarpSize = cutlass::gemm::warp::WarpSize<arch::OpClassSimt>::value;
                                                                               ~~^
./cutlass/include/cutlass/epilogue/threadblock/default_epilogue_simt.h:367:5: error: use of class template 'OutputTileThreadMap' requires template arguments
    OutputTileThreadMap,
    ^
./cutlass/include/cutlass/epilogue/threadblock/output_tile_thread_map.h:134:8: note: template is declared here
struct OutputTileThreadMap : public OutputTileThreadMapHelpers<Iterations_, Delta_> {
       ^
In file included from cutlass/include/cutlass/epilogue/threadblock/default_epilogue_simt.cu:1:
./cutlass/include/cutlass/epilogue/threadblock/default_epilogue_simt.h:391:5: error: use of class template 'OutputTileThreadMap' requires template arguments
    OutputTileThreadMap,
    ^
./cutlass/include/cutlass/epilogue/threadblock/output_tile_thread_map.h:134:8: note: template is declared here
struct OutputTileThreadMap : public OutputTileThreadMapHelpers<Iterations_, Delta_> {
       ^
In file included from cutlass/include/cutlass/epilogue/threadblock/default_epilogue_simt.cu:1:
./cutlass/include/cutlass/epilogue/threadblock/default_epilogue_simt.h:405:5: error: unknown type name 'OutputTileIterator'; did you mean 'WarpTileIterator'?
    OutputTileIterator,
    ^
./cutlass/include/cutlass/epilogue/threadblock/default_epilogue_simt.h:380:9: note: 'WarpTileIterator' declared here
  using WarpTileIterator = cutlass::epilogue::warp::TileIteratorSimtDirect2dConv<
        ^
./cutlass/include/cutlass/epilogue/threadblock/default_epilogue_simt.h:408:5: error: use of class template 'SharedLoadIterator' requires template arguments
    SharedLoadIterator,
    ^
./cutlass/include/cutlass/epilogue/threadblock/shared_load_iterator.h:67:7: note: template is declared here
class SharedLoadIterator {
      ^
```
2022-12-21 11:40:20 -05:00
3f2bb17722 minor chagnes (#730)
Co-authored-by: Haicheng Wu <haichengw@nvidia.com>
2022-12-10 14:44:53 -05:00
38193d76e3 Updates for stream-k (#728)
Co-authored-by: Aniket Shivam <ashivam@nvidia.com>
2022-12-08 23:48:10 -05:00
1d7772f218 Add missing #include directive (#727) 2022-12-08 18:58:31 -05:00
df81d847d7 Make Python interface work for non-SM80 targets (#726)
* Make Python interface work for non-SM80 targets

* Remove line in README
2022-12-07 21:53:33 -05:00
d6117ca362 Relax stream K gemm alignment constraints (#717)
* Relax stream K gemm alignment constraints

The current alignment requirements are too strict. Make them identical
to the checks for the regular universal gemm.

* Revert "Relax stream K gemm alignment constraints"

This reverts commit 31e80a250e.

* Relax stream K gemm alignment constraints

The current alignment requirements are too strict. Make them identical
to the checks for the regular universal gemm.

Co-authored-by: Haicheng Wu <haichengw@nvidia.com>
2022-12-07 11:17:49 -05:00
9c0518608e Fix typos in conv problem sizes (#720)
* Fix typos in conv problem sizes

* Typos
2022-12-05 15:54:58 -05:00
9f1f37aa21 misc (#719)
* misc

* minor

Co-authored-by: Haicheng Wu <haichengw@nvidia.com>
2022-12-05 12:07:20 -05:00
84213b0b8e fix: make arch.h self contained (#714) 2022-12-01 19:25:48 -05:00
8567b87d65 Update quickstart.md (#704)
* Update quickstart.md

* Update doxygen_mainpage.md

* Update doxygen_mainpage.md

* Update terminology.md
2022-11-29 21:43:03 -05:00
c975e2ccbb releaase 2.11 (#703) 2022-11-19 09:02:15 -05:00
3c90f6aea6 add #pragma once for header file in example 42 (#698) 2022-11-15 22:50:24 -05:00
06eb90cc0d Fix identity sigmoid activation (#659)
* activation support Identity

* fix Sigmoid activation operator() with CUTLASS_HOST_DEVICE
2022-11-09 14:42:23 -05:00
168ea8b0e1 ensure singleton::get thread safe construct instance (#658)
* ensure singleton::get thread safe construct instance

* fix singleton return reference

Co-authored-by: xuweiqi <xuweiqi117@gmail.com>
2022-11-08 21:44:32 -05:00
012c62c748 bug fixes and enharcement to gemm reductionK fusion (#682)
* add two missing files

* fix bunch of bugs of gemm-reducek fusion and add a device interface

* small changes

Co-authored-by: Haicheng Wu <haichengw@nvidia.com>
2022-11-03 11:07:50 -04:00
FZC
cc85b64cf6 fix typo (#677) 2022-11-01 14:07:33 -04:00
1b4e24470a Example 43 - DualGemm (#670)
* Ex50 wip

* IS_PROFILING mode

* MultiStage2 - but is slower

* Add SwiGLU

* Support SplitKSerial reduction
Support not storing D0/D1
Cleanup code

* Option to disable bias

* Renumber example

* Fix build

* Remove references to pb_size_0 / pb_size_1

* Add support for bf16 inputs with float accum

* small changes

Co-authored-by: danthe3rd <danthe3rd>
Co-authored-by: Haicheng Wu <haichengw@nvidia.com>
2022-10-26 14:04:42 -04:00
8c1bf9b784 Bump CUTLASS Python container version (#672)
* Update example 40 README

* Update CUTLASS Python README
2022-10-22 21:09:39 -04:00
7d0dd6706e Remove excessive includes from examples/41_multi_head_attention (#669)
The rationale behind this change is explained in #563
2022-10-21 22:23:15 -04:00
9b47403b2d Add missing CUTLASS_HOST_DEVICE (#671) 2022-10-21 22:20:38 -04:00
4db6a6140e ex42: Fused MHA imported from xFormers (#662)
* ex42: Fused MHA imported from xFormers

* Remove std:: references

* Support K>128 in the example

* Support causal option

* Support different head size for V, and different seqlength for KV

* Update FLOPS counter

* Remove bit_cast

* fix build: Replace M_LOG2E

* Add doc

* Revert "Remove bit_cast"

This reverts commit 9662fa86bb.

* Explicit casts to int32_t for windows build

Co-authored-by: danthe3rd <danthe3rd>
2022-10-17 10:49:33 -04:00
3bf95e90c2 Update labeler.yml 2022-10-13 08:03:28 -04:00
75fed7493e Update labeler.yml 2022-10-13 08:01:21 -04:00
98b73fc95d Update labeler.yml 2022-10-13 07:55:33 -04:00
4990e3686d Update labeler.yml 2022-10-13 07:52:38 -04:00
4b7365388c Update labeler.yml 2022-10-13 07:32:55 -04:00
0d8405588d Update labeler.yml 2022-10-12 15:32:38 -04:00
cb539dab78 Correct typos in comments (#639)
* Correct typos in comments

Correct comments in code on type of generated distribution. Improve Gaussian RNG to take advantage of Box Muller method

* Inline Box Muller

Added inline function for the Box Muller algorithm and updated code comments to be more concise

* Update tensor_fill.h

* Update tensor_fill.h

* small changes to pass tests

Co-authored-by: Haicheng Wu <haichengw@nvidia.com>
2022-09-30 22:51:30 -04:00
dadc881a96 Bug fix for gemm broadcast (#650)
* gemm_universal_with_broadcast, +2 sources.

* Revert "gemm_universal_with_broadcast, +2 sources."

This reverts commit fb063251f2.

* gemm broadcast bug fix
2022-09-30 10:00:38 -04:00
f3eea3a4d7 Create labeler.yml 2022-09-29 15:08:44 -04:00
cd37e82492 change unused class member to local var (#646) 2022-09-28 23:52:35 -04:00
48a9ea223a Fix release version in the citation (#638) 2022-09-22 10:58:45 -04:00
7a458f00a6 fix(permute.h): incorrect comment in Tensor5DPermute20314 (#637)
* fix(permute.h): incorrect comment in `Tensor5DPermute20314`

* typo in usage in example 39
2022-09-22 09:21:13 -04:00
97bff52e8c add two missing files (#636)
Co-authored-by: Haicheng Wu <haichengw@nvidia.com>
2022-09-21 15:42:42 -04:00
9f2e3faa69 fix call of GELU_Taylor in LinearCombinationGeneric (#634) 2022-09-20 21:00:55 -04:00
a821280dc7 Gemm broadcast (#632)
* gemm_universal_with_broadcast, +2 sources.

* Revert "gemm_universal_with_broadcast, +2 sources."

This reverts commit fb063251f2.

* gemm_universal_with_broadcast separated version.

* Update copyright banner.

* update banner
2022-09-20 10:37:12 -04:00
f73374a1eb fix:comment typo in example 23 (#633) 2022-09-19 09:54:14 -04:00
faab7536fc add comment (#628) 2022-09-17 21:40:30 -04:00
fc9ebc645b CUTLASS 2.10 bug fixes and minor updates. (#626) 2022-09-15 16:20:33 -04:00
2cc2c7ba1f Add set_k_partition function (#624)
A member function set_k_partition is required for the instatiation of cutlass::gemm::kernel::Gemm, even though SplitKSerial is false
2022-09-13 22:34:20 -04:00
50ceed7154 Minor README fix (#623)
* minor fix

* Minor fix
2022-09-12 22:40:25 -04:00
e773429f7e CUTLASS 2.10 updates (#622)
Co-authored-by: Aniket Shivam <ashivam@nvidia.com>
2022-09-12 21:26:30 -04:00
beae168f90 fix broken link (#620)
Co-authored-by: yuzhai <yuzhai@nvidia.com>
2022-09-06 16:32:44 -04:00
f29d8f7ca9 Include vector in base_grouped.h (#618) 2022-09-06 13:21:23 -04:00
b1d3f9b2fd upstream internal updates (#616)
Co-authored-by: yuzhai <yuzhai@nvidia.com>
2022-09-04 23:05:09 -04:00
b72cbf957d CUTLASS 2.10 (#615)
Co-authored-by: Aniket Shivam <ashivam@nvidia.com>
2022-09-03 18:48:46 -04:00
ca23ff7924 Fixed typo in class name (#608) 2022-08-29 20:51:52 -04:00
1c3d400b14 Added value_type trait to complex to make it an easier drop-in replacement for std::complex. (#607) 2022-08-28 01:12:40 -04:00
abafbf2afd Missing comma in trmm header (#604) 2022-08-25 16:07:33 -04:00
536b20763e Fixed typo in profiler README (#603) 2022-08-24 21:55:13 -04:00
497b499d9d Add residual support for shmem staging iterator used in back-to-back GEMM fusion. This allows support of problem_size_0_n that is not multiple of 32. (#590)
Co-authored-by: Haicheng Wu <haichengw@nvidia.com>
2022-08-15 11:19:24 -04:00
e66bfcb1f8 Fix for #596 (typo in example 03) (#597)
* [examples] Fix typos in SYRK and TRMM examples

* Fix typo in example 03
2022-08-09 09:58:36 -04:00
1617685a77 fix: fix types in example 06 (#587) 2022-07-29 12:46:06 -04:00
25ebf15d02 Ensure all arch::Mma specializations have ElementC set (#576)
Co-authored-by: danthe3rd <danthe3rd@users.noreply.github.com>
2022-07-22 23:53:03 -04:00
5d05808072 fix gather example (#574) 2022-07-19 16:18:17 -04:00
0b8cacd6f1 Remove redundant <fstream> includes (#563)
* Remove redundant <fstream> includes

* Fix fstream in examples/

* Fix <fstream> in test/

* Use consistent order for <fstream> (always after <iostream>)

* Remove an unneeded include in a file where std::ofstream usage is commented out

Co-authored-by: Ivan Komarov <dfyz@yandex-team.ru>
2022-07-19 15:23:54 -04:00
e7a61c761a fix race condition when h < stride_h or w < stride_w (#562)
Co-authored-by: Haicheng Wu <haichengw@nvidia.com>
2022-07-12 16:37:08 -04:00
fb379eaa5b epilogue leaky relu support ScaleType (#564)
Co-authored-by: xuweiqi <xuweiqi117@gmail.com>
2022-07-11 17:30:55 -04:00
8a766804ad Fix doc in testbed_gemm_with_broadcast (#559) 2022-07-07 09:56:16 -04:00
1eb6355182 [activation] tanh (#550)
Co-authored-by: Bing Xu <bingxu@fb.com>
2022-07-02 08:00:45 -04:00
04a9777b87 Softmax (#546)
* add test layernorm g-mem version

* Delete include/configure directory

* Delete examples/test_layernorm directory

* Update gemm_with_softmax.h

* Update gemm_softmax.cu

* Update linear_combination.h

* Update fast_math.h

* remove redundant vars

Co-authored-by: yujia.zhai <yujia.zhai@bytedance.com>
Co-authored-by: yuzhai <yuzhai@nvidia.com>
2022-07-02 01:19:18 -04:00
e45e773436 Update linear_combination_generic.h (#472)
add `skip_elementwise_` to support serial splitk in linear_combination_generic.h`
2022-06-28 07:29:38 -04:00
dae6b6893b Update CHANGELOG.md 2022-06-27 23:30:49 -04:00
ba18ea9c32 Update README.md 2022-06-27 23:25:26 -04:00
9ab9110168 add leaky relu (#542)
Authored-by: Haicheng Wu <haichengw@nvidia.com>
2022-06-26 10:07:50 -04:00
e5d4669f16 Update CHANGELOG.md (#543) 2022-06-25 13:23:49 -04:00
94f01f19d5 Add implicit gemm perf
plot from @manishucsd, presented in gtc'22 cutlass talk
2022-06-23 22:47:11 -04:00
fa56763c25 Fix occupancy calculation for grouped GEMM (#532) 2022-06-18 19:53:59 -04:00
25e26a6e51 fix bugs in linear_combination_generic.h missing include cutlass/epilogue/thread/scale_type.h (#531) 2022-06-17 23:35:14 -04:00
f248e9bdb4 Create CITATION.cff
Add initial CITATION.cff
2022-06-07 21:25:16 -04:00
dceefe4f64 Increment stride correctly in warp iterator. (#516)
Co-authored-by: peisun1115 <peis@google.com>
2022-06-06 12:33:36 -04:00
c3881d097e Fix a comment about LDSM layout. (#514)
Co-authored-by: peisun1115 <peis@google.com>
2022-06-04 23:04:00 -04:00
a29dfb1c63 Fix a bug to increment stride tile correctly (#503)
* Fix a bug to increment stride tile correctly

* Update regular_tile_access_iterator_tensor_op.h

Co-authored-by: peisun1115 <peis@google.com>
Co-authored-by: Haicheng Wu <57973641+hwu36@users.noreply.github.com>
2022-06-03 22:54:52 -04:00
0abaac84ea [examples] Fix typos in SYRK and TRMM examples (#507) 2022-06-03 22:52:41 -04:00
858c735856 Update gather_scatter_fusion.cu
Correct the reference code in gather/scatter example to put bias add in the correct place.
2022-05-18 13:15:25 -04:00
d6f58b2d14 Update functionality.md 2022-05-11 09:34:24 -04:00
c4cf0dad82 Fix init-self compiler warnings (#493)
Fix a few errors caused by trying to initialize a class member
with itself. These errors can turn into errors if you compile
with `-Winit-self`.
2022-05-11 00:35:28 -04:00
57551902d0 Update functionality.md
add some explanations to the functionality table.
2022-05-11 00:01:19 -04:00
1604ebaf10 Update generator.py
stop generating analytical conv kernels to reduce kernel number
2022-05-08 21:47:15 -04:00
6023038bae add verification of the reduction tensor (#489)
Co-authored-by: Haicheng Wu <haichengw@nvidia.com>
2022-05-06 10:24:51 -07:00
ddd8f9cf41 update float < int32_t * 4 (#488)
Co-authored-by: 赵俊涛 <zhaojuntao@zhaojuntaos-MacBook-Pro.local>
2022-05-04 13:36:05 -04:00
ec2b4fd85d b2b bias vector support (#482)
* b2b bias vector support

* add files

Co-authored-by: Haicheng Wu <haichengw@nvidia.com>
2022-04-30 04:16:15 -07:00
86ce09aed1 2.9 fixes for nvrtc (#480)
* Use platform::is_same instead of std::is_same

* Don't hide cuComplex include from nvrtc

* Typo fixed

* Remove comment rename
2022-04-29 09:06:52 -04:00
21c1fa3849 add .github (#479)
Co-authored-by: Haicheng Wu <haichengw@nvidia.com>
2022-04-28 12:36:59 -07:00
8c339ac039 Fix compilation in clang (#478)
- adds missing commas
- adjusts misaligned usage of CUTLASS_DEVICE between
  template declaration and specializations

Signed-off-by: Janusz Lisiecki <jlisiecki@nvidia.com>
2022-04-28 14:22:06 -04:00
e49f690fd7 Update linear_combination_generic.h 2022-04-28 14:04:53 -04:00
96dad61a75 Update CHANGELOG.md 2022-04-28 10:52:10 -04:00
cc2ea4c3fc Update README.md 2022-04-28 10:50:11 -04:00
a0de301283 Used relative paths for includes (#477) 2022-04-27 12:04:23 -07:00
319a389f42 Update CMakeLists.txt (#473)
* Update CMakeLists.txt

Add 128bit int support if using nvc++ to solve #310 

@jeffhammond, would you please give it a try?

* Update CMakeLists.txt

correct copy paste error
2022-04-27 07:02:26 -07:00
71def2f084 Use platform:: instead of std::abs and std::conditional (#452)
* Fixed template struct/class mismatch

* Use platform implementation instead of std::abs and std::conditional during nvrtc compilation

* Use platform implementation instead of std::abs and std::conditional during nvrtc compilation

* Revert absolute_value() usage
2022-04-25 14:40:22 -04:00
70f3ba57f5 Fix typo in shared memory layout description (#471) 2022-04-24 18:32:13 -04:00
dd77fadc70 Remove redundant offset def and init in shared_load_iterator.h (#456)
Signed-off-by: Fujun Han <fujun.han@iluvatar.ai>
2022-04-24 16:31:00 -04:00
be4578d517 Fixed template struct/class mismatch (#453) 2022-04-24 16:30:21 -04:00
d7b499deff Fix CUDA_PERROR_EXIT and print failing expression (#446)
`CUDA_PERROR_EXIT ` can lead to incorrect usage (see e.g. [this description](https://www.cs.technion.ac.il/users/yechiel/c++-faq/macros-with-if.html)) because it contains an incomplete `if` expression. Consider:

```
if (condition)
    CUDA_PERROR_EXIT(cudaFree(x))
else
    free(x);
```

The author of the code forgot to add a semicolon after the macro. In that case, the `else` will bind to the `if` inside the macro definition, leading to code that the author did not intend or expect. It the author does use a semicolon, the code will not compile, which is awkward.

The change adds a `do while` around the `if`, which always requires a semicolon.

This PR also adds the text of the failing expression to the printed error message.
2022-04-24 16:29:43 -04:00
310ed81ac3 fix description in example 12. (#444)
Co-authored-by: Exusial <Exusial>
2022-04-24 16:29:06 -04:00
4c0d6e1eb4 [BUGFIX]: Force unroll a loop that doesn't have compilation constant (#441)
loop times is dangerous.

Signed-off-by: Peter Han <fujun.han@iluvatar.ai>
2022-04-24 16:28:32 -04:00
167ac54c65 Fix link to Python example (#469) 2022-04-23 15:37:38 -04:00
12f4108ac2 CUTLASS 2.9 (#468) 2022-04-23 15:02:38 -04:00
dd571f0edb [style] fix code indentation (#449)
* [docs] fix typo in media/docs/layout.md

* [docs] fix comment error

* fix typo in include/cutlass/arch/simd_61.h

* fix stride comment errors in TensorLayout

* fix indentation
2022-04-03 21:13:17 -04:00
6d0d265047 Update PUBLICATIONS.md (#447) 2022-04-03 21:03:28 -04:00
f11fa975a5 Update PUBLICATIONS.md
@tsuki
2022-03-23 21:04:43 -04:00
0e71d9b450 Transposed conv2d and wgrad split k examples (#413)
* add split k wgrad example

* wgrad done

* begin transposed conv2d example

* update transposed conv2d example and add ref check

* update doc for conv2d transpose example

* add license

* add wgrad doc

* more clarification on GEMM output type

* typo fix

* clean up indent

* address comments

* rename example numbers to 34 and 35

* GEMM -> Implicit GEMM

* Revert "rename example numbers to 34 and 35"

This reverts commit 551a808c22.

* transposed_conv2d is 34

* add compiler and device version check to exit gracefully

Co-authored-by: Haicheng Wu <haichengw@nvidia.com>
2022-03-23 14:52:54 -04:00
eb0d4c9213 [library] pass pointer of arguments to get_host_workspace_size() in gemm_universal() (#412)
Otherwise GemmUniversalOperation::get_host_workspace_size() will fail on SegmentFault.
2022-03-22 12:36:34 -04:00
bc45e2c023 fixed datatype error of numeric_limit for uint1b_t (#419)
Co-authored-by: Haojin Yang <haojin.yang@.hpi.uni-potsdam.de>
2022-03-22 12:30:30 -04:00
095cbba57c Example 23 - Passing correct alpha and beta values with --parallel-split-k (#424)
When split-k is enabled, we should set alpha to 1 and beta to 0 for the
split-k gemm kernel.

The fix was from hwu36. I only did fixed some minor typos along with his
fix.
2022-03-22 12:27:34 -04:00
8f1fe7a132 Fix separate compilation -dc (#433)
* Fix separate compilation `-dc`

- when cutlass is included in multiple compilation units
  compiled with `-dc` OOB_NAN_F16x8 device constant is
  instantiated multiple times causing
  Multiple definition of '_ZN7cutlass4arch13OOB_NAN_F16x8E' error
  This PR makes this variable a local constant as it is not
  modified during runtime

Signed-off-by: Janusz Lisiecki <jlisiecki@nvidia.com>

* Fix

Signed-off-by: Janusz Lisiecki <jlisiecki@nvidia.com>

* Test GH

Signed-off-by: Janusz Lisiecki <jlisiecki@nvidia.com>

* Revert test GH

Signed-off-by: Janusz Lisiecki <jlisiecki@nvidia.com>
2022-03-22 12:21:18 -04:00
3ab1eacf09 Fix typo in profiler examples (#437) 2022-03-21 12:00:13 -04:00
cd39c75e25 Fix typo in docs, code comments (#429)
* [docs] fix typo in media/docs/layout.md

* [docs] fix comment error

* fix typo in include/cutlass/arch/simd_61.h

* fix stride comment errors in TensorLayout
2022-03-15 21:54:36 -04:00
b2e1e97cb1 Update PUBLICATIONS.md
ACM Trans on Graphics from nv research.
2022-03-01 22:37:18 -05:00
96a11a1ef3 Removed trivial copy constructors on parameter classes to enable devi… (#366)
* Removed trivial copy constructors on parameter classes to enable device-side launch of CUTLASS kernels

* Added SFINAE to the `TensorRef(NonConstTensorRef const&)` constructor to avoid making it a copy-constructor for device code

* std => platform

* fix affine2

* really fix affine2

Co-authored-by: Haicheng Wu <haichengw@nvidia.com>
2022-02-28 21:34:02 -05:00
e96f00586c Make cutlass::gemm::device::GemmArray usable (#295)
* Fix the build of cutlass/gemm/device/gemm_array.h and add a demo for GemmArray

* Add a reference to GemmArray to the docs

Co-authored-by: Ivan Komarov <dfyz@yandex-team.ru>
2022-02-17 20:01:05 -05:00
3cfa5db2a2 Actually use float accumulation in gemm_f16t_f16t_f16t_wmma_tensor_op… (#407)
* Actually use float accumulation in gemm_f16t_f16t_f16t_wmma_tensor_op_f32_sm70.cu

As title

* Update gemm_f16t_f16t_f16t_wmma_tensor_op_f32_sm70.cu

change the missing one

Co-authored-by: Haicheng Wu <57973641+hwu36@users.noreply.github.com>
2022-02-16 09:53:21 -05:00
1db6971a8d Remove unused gemm_k_iterations in GemmKernel::Params (#406)
Otherwise we get gemm_k_iterations is uninitialized warnings.
2022-02-16 09:52:45 -05:00
b954127297 Update PUBLICATIONS.md
@jackkosaian
2022-02-14 16:54:32 -05:00
d0d941efc7 [hardswish] correct implmentation (#403)
* [hardswish] correct implmentation

* seems working

* hardswish fp32/fp16x2 optimization

* [relu] half2 support

* add relu0; add multiply_add_relu0;

* cleanup

Co-authored-by: Bing Xu <bingxu@fb.com>
Co-authored-by: Haicheng Wu <haichengw@nvidia.com>
2022-02-09 14:28:53 -05:00
8a951b2940 Enable convolution with fused epilogue for Volta Tensor Cores (#402)
* Enabled convolution with epilogue fusion for Volta Tensor Cores.

* Compilation fixes

* Disabled testing Volta on Ampere architectures.
2022-01-30 23:24:50 -05:00
1e4703cbab Support parallel split K mode for porfiling (#277)
* Support parallel split K mode for porfiling

Signed-off-by: Peter Han <fujun.han@iluvatar.ai>

* Parallel Split K support

  1. find gemm kernel by preference key
  2. switch m n for redution kernel

Signed-off-by: Peter Han <fujun.han@iluvatar.ai>

* parallel splitk for fp16 gemm

* add one missing file

Co-authored-by: Haicheng Wu <haichengw@nvidia.com>
2022-01-27 10:37:37 -05:00
c3353add63 Merge pull request #388 from depaulmillz/fix/headersonly
Fix utils include not being installed in header only
2022-01-26 14:22:51 -06:00
ac8825b941 Minor fix to change from LIBRARY_INIT to LIBRARY 2022-01-26 15:17:46 -05:00
8fd94806e5 Update PUBLICATIONS.md
add mlsys 2022 paper.
2022-01-17 00:08:18 -05:00
d7c9cbf0b9 Fix typo in scripts/library.py (wrong data size for u8) (#393) 2022-01-07 13:29:56 -05:00
c2ee13a0fe Add epilogue functor for residual block fusion (#391)
* Add epilogue functor for residual block fusion

* Do not run split-k tests when ActivationOp is not Identity

* explain TestSplitK param

* return early
2021-12-29 22:53:40 -05:00
f78994bb40 add the missing pieces (#392)
Co-authored-by: Haicheng Wu <haichengw@nvidia.com>
2021-12-25 04:29:54 -08:00
dceabd4c5a Support half precision sigmoid activation (#378)
* Support half precision sigmoid activation

* introduce a vectorized variant using fast_tanh

* move the math to fast_math.h

* fixed compile

* .raw() -> .to_half()

Co-authored-by: Haicheng Wu <haichengw@nvidia.com>
2021-12-22 14:45:06 -05:00
86fa1dc30b Fix utils include not being installed in header only 2021-12-21 12:10:26 -05:00
288af365db Added missing synchronization to avoid WAR hazards between tiles. (#386) 2021-12-20 08:34:08 -08:00
0dc3ba60b3 Refactor GELU and Sigmoid epilogue to use a common template (and add SiLu, Hardswish epilogue) (#379)
* Support half precision sigmoid activation

* introduce a vectorized variant using fast_tanh

* refactored sigmoid using the new interface

* refactored gelu

* add silu activation

* add hardswish

* remove sigmoid for now

* add description to silu and hardswish, and other doc update

* Do not ignore Round

* use constant N

* Set isHeavy = true in sigmoid and silu epilogue
2021-12-18 14:58:15 -05:00
ec4f7e5194 Updates to fused epilogue (#383)
* Enhancements and fixes to fused GEMM and Convolution epilogue.
* Need to explicitly list cudart as unit test library dependency.
2021-12-17 16:04:43 -05:00
4e666e1dfd Updated README and added issue templates. (#382) 2021-12-17 09:26:20 -05:00
3799e12f25 Merge pull request #381 from Peter9606/update-makefile-version
Update project version to 2.8.0 in CMakeLists.txt
2021-12-16 21:54:57 -05:00
fc3bc85db8 Update project version to 2.8.0 in CMakeLists.txt
Signed-off-by: Peter Han <fujun.han@iluvatar.ai>
2021-12-17 02:23:31 +00:00
49c0a58d50 Set theme jekyll-theme-minimal 2021-12-15 14:51:24 -05:00
5fe09c2d67 Updated GEMM performance plot with CUTLASS 2.8 compiled with CUDA 11.5 Toolkit (#375)
Updated GEMM performance plot with CUTLASS 2.8 compiled using CUDA 11.5 Toolkit.

GPUs under test:

    NVIDIA A100
    NVIDIA A2
    NVIDIA TitanV
    NVIDIA GeForce 2080 Ti
2021-12-06 14:21:33 -05:00
6b69c79ac3 Fixed contributor formatting. (#365) 2021-11-22 11:30:53 -08:00
62e438f450 Listed Matthew Nicely as the CUTLASS product manager.. (#364) 2021-11-19 17:51:21 -08:00
808c25337a CUTLASS 2.8 (#363)
CUTLASS 2.8
2021-11-19 13:26:35 -08:00
6fc5008803 Update quickstart.md
fix a broken link
2021-11-11 09:53:46 -05:00
a3bcc6981d Merge pull request #331 from reed-lau/feature/fix-wmma-shape-typo
fix wmma shape typo
2021-09-28 10:20:29 -04:00
3b28642801 fix wmma shape typo 2021-09-28 19:04:09 +08:00
538592dea4 example 23 gemm operand reduction fusion (#325) 2021-09-20 13:34:47 -07:00
2e07c4cc2f CUTLASS 2.7 (#318)
CUTLASS 2.7

Mainloop fusion for GEMM: summation over A or B
Strided DGRAD (optimized iterators)
Half-precision GELU_taylor activation functions
Use these when accumulation and epilogue compute types are all cutlass::half_t
Tuning and bug fixes to fused GEMM + GEMM example
Support for smaller than 128b aligned Convolutions: see examples
Caching of results to accelerate Convolution unit tests
Can be enabled or disabled by running cmake .. -DCUTLASS_TEST_ENABLE_CACHED_RESULTS=OFF
Corrections and bug fixes reported by the CUTLASS community
Thank you for filing these issues!

authored-by: Haicheng Wu haichengw@nvidia.com, Manish Gupta manigupta@nvidia.com, Dustyn Blasig dblasig@nvidia.com, Andrew Kerr akerr@nvidia.com
2021-09-20 11:02:22 -07:00
9ac255863f Merge pull request #246 from mengchihe/master
support unalignment input for conv2d fprop stage=2 Fix for issue #242
2021-09-08 11:40:53 -04:00
59e2aa505a refine the implementation 2021-09-08 13:14:08 +00:00
4e8af93da1 Merge remote-tracking branch 'origin/master' into small_alignment 2021-09-07 20:39:38 +00:00
6c2f8f2fb8 CUTLASS 2.6.1 - functional and performance enhancements to strided DGRAD, fixes, and tuning
* cutlass 2.6 update

* remove debug prints

* cutlass 2.6.1 (minor update)

* Updated CHANGELOG.

* Minor edit to readme to indicate patch version.

* Minor edit to readme.

Co-authored-by:  Haicheng Wu <haichengw@nvidia.com>, Andrew Kerr <akerr@nvidia.com>
2021-09-03 10:26:15 -07:00
598e35401c Merge remote-tracking branch 'origin/master' into small_alignment 2021-08-16 07:49:08 -07:00
a01feb93d9 Merge pull request #308 from dongxiao92/patch-1
fix typo in doc
2021-08-08 11:54:42 -07:00
d36f331b44 fix typo in doc
fix typo
2021-08-08 16:44:22 +08:00
69abafb85a Merge pull request #306 from NVIDIA/fix-profiler-cmd-doc
Fix profiler cmd doc
2021-07-30 14:36:54 -04:00
68a078fbbf cleanup 2021-07-30 11:27:21 -07:00
10709dbb64 clean profiler cmd and doc 2021-07-30 11:02:17 -07:00
1227351079 Merge pull request #305 from NVIDIA/fix_epilogue_spill
fix epilogue register spill
2021-07-29 14:30:11 -07:00
a77c658439 fix epilogue register spill 2021-07-29 14:25:48 -07:00
4516b833ce Merge pull request #303 from Peter9606/doc_typo
Doc typo
2021-07-28 20:49:06 -04:00
64dd1e1915 Doc typo
Signed-off-by: Peter Han <fujun.han@iluvatar.ai>
2021-07-29 08:45:59 +08:00
1ac4559d12 Cutlass 2.6 Update 1 (#301)
* cutlass 2.6 update

* remove debug prints
2021-07-27 17:58:30 -07:00
e5d51840e8 CUTLASS 2.6 (#298)
CUTLASS 2.6
2021-07-23 00:40:53 -04:00
6c29fe20ba Merge pull request #285 from tjingrant/patch-1
Typo Fixes
2021-07-05 22:51:19 -04:00
e3c56b0d6b Update predicated_tile_iterator.h 2021-07-05 12:11:53 -04:00
4647c57243 Update predicated_tile_iterator.h 2021-07-05 12:06:41 -04:00
856d4db3fb Update basic_gemm.cu
fix the matrix malloc size
2021-06-15 09:08:36 -04:00
6a1064093f Merge pull request #274 from mani-ananth/master
Some pending Bug fixes
2021-06-02 13:17:39 -04:00
c5f1ef4dff update contributors 2021-06-02 10:11:42 -07:00
47ebfccbec bug fixes 2021-06-02 10:08:25 -07:00
ad9486684f Merge pull request #272 from BernardoCovas/master
Bug in reference conv3d
2021-05-28 17:18:27 -04:00
1d8372a8e2 fix typo in reference conv3d 2021-05-28 21:06:59 +01:00
9cb7d63424 Merge pull request #266 from mani-ananth/master
Fixes for public issue #265
2021-05-19 15:15:22 -04:00
da2f110906 Fixes for public issue #265 2021-05-19 10:16:52 -07:00
b68113f5be Merge pull request #264 from zheng95z/patch-3
Adds `NoBetaScaling` for `LinearCombination`
2021-05-17 10:03:30 -04:00
a68d7cd6f1 Adds NoBetaScaling for LinearCombination 2021-05-12 22:23:55 +08:00
38e8b29f56 Merge pull request #259 from hzfan/ignore_pr
Add gitignore
2021-05-10 20:07:53 -04:00
ee7349c94f fix 2021-05-10 16:39:04 +08:00
8cdd4293d4 add gitignore 2021-05-10 16:37:59 +08:00
f58b843951 Merge pull request #239 from KeDengMS/kedeng/gelu
Fixes to Gelu for half and fusion
2021-05-08 12:51:42 -04:00
5fc142296f Merge pull request #237 from Peter9606/issue_236_typo
Typo fix issue#236
2021-05-08 07:51:19 -04:00
233d69aa6d Merge pull request #235 from Peter9606/issue_233_tranpose_update
tranpose.h update based on issue#233
2021-05-07 07:14:30 -04:00
9840d25269 Merge pull request #256 from zheng95z/patch-2
Fixes some typos in utilities.md
2021-05-06 11:02:49 -04:00
b878c96421 Fixes some typos in utilities.md 2021-05-06 22:37:37 +08:00
8f8a80cad5 Merge pull request #251 from zheng95z/patch-1
add a missing 'device_memory::' before a function
2021-04-25 22:09:44 -04:00
a8f6f8eb07 add a missing 'device_memory::' before a function 2021-04-25 20:05:39 +08:00
f4b0a33633 add unit test for non int4 load 2021-04-23 14:33:46 +08:00
7c783adf53 Merge pull request #247 from xue-fc/patch-1
fix a wrong description
2021-04-22 09:27:40 -04:00
4000df9567 fix a wrong description 2021-04-22 20:28:28 +08:00
bb35a3ba6f support setting load granularity for conv2d fprop 2021-04-22 15:20:57 +08:00
7ec3a87f22 support unalignment input for conv2d fprop stage=2 Fix for issue #242 2021-04-21 14:40:05 +08:00
0b74c8f473 Address CR 2021-04-19 23:36:06 +00:00
83036ed646 More clean up 2021-04-18 04:29:20 +00:00
b7e43f5eb9 Clean up 2021-04-18 04:24:25 +00:00
5c62d892fa Add test 2021-04-18 04:09:34 +00:00
41a31b404b Fixes to Gelu for half and fusion 2021-04-17 22:10:19 +00:00
7320aee17d Typo fix issue#236
Signed-off-by: Peter Han <fujun.han@iluvatar.ai>
2021-04-15 15:08:35 +08:00
2142a05d9d tranpose.h update based on issue#233
1. Add 'pragma once' preprocess directive
 2. Replace prmt PTX with __byte_perm intrinsic

Signed-off-by: Peter Han <fujun.han@iluvatar.ai>
2021-04-14 19:58:00 +08:00
c77a524459 Merge pull request #230 from mani-ananth/master
Fix for issue #221
2021-04-09 14:45:55 -04:00
fac6680f31 Merge branch 'master' of github.com:NVIDIA/cutlass 2021-04-09 11:36:31 -07:00
08993707da fixing functional bug in fused epilogue 2021-04-09 11:36:03 -07:00
c805593ebe Merge pull request #228 from mani-ananth/master
Fix for issue#224 and issue#225
2021-04-08 10:08:13 -04:00
26556d7206 fix a broken sparse gemm example. found by the community. 2021-04-07 13:32:55 -07:00
4839b6cb61 add 2stage fprop 3d into default file 2021-04-07 13:29:32 -07:00
d97214987a Merge pull request #220 from Peter9606/wrong-stride-array-definition
Bugfix: typo, make reduction device cases passed
2021-04-02 08:43:52 -04:00
b0bbc6d548 Merge pull request #219 from mani-ananth/master
Fix for issue #211
2021-04-02 08:42:09 -04:00
7074047a54 Bugfix: typo, make reduction device cases passed
Signed-off-by: Peter Han <fujun.han@iluvatar.ai>
2021-04-02 09:35:23 +08:00
75a4737cfe Fix for public issue #211
- Add a slice-K tile size to the profiler
- fix num warps calculations in implicit gemm header
2021-04-01 14:42:00 -07:00
8a3e4b8d02 Merge pull request #214 from Peter9606/separate-stream-error
Bugfix: memsetAsync uses wrong default stream
2021-03-24 12:09:01 -04:00
6a6b4028bd Revert wrong fix of params.update in GemmUniversalBase
Signed-off-by: Peter Han <fujun.han@iluvatar.ai>
2021-03-23 23:20:40 +08:00
92393b2676 Bugfix: memsetAsync uses wrong default stream
Signed-off-by: Peter Han <fujun.han@iluvatar.ai>
2021-03-23 21:11:42 +08:00
50bf00e5f2 Merge pull request #193 from Peter9606/public_shape_type_from_Mma_HFMA2
HFMA2 Convolutions for SM60 onwards
2021-03-05 21:38:59 -05:00
4cd004ead1 fix test name to optimized and instance large tile sizes to speed unit tests 2021-03-05 13:32:36 -08:00
6c4539e372 Make arch tag of test cases more precisely to SM60
Signed-off-by: Peter Han <fujun.han@iluvatar.ai>
2021-03-05 10:53:26 +08:00
a3639ab1a0 Append fp16 test case to verify Mma_HFMA2
Signed-off-by: Peter Han <fujun.han@iluvatar.ai>
2021-03-04 18:17:57 +08:00
169181f30f Make Shape public from Mma_HFMA2.
Signed-off-by: Peter Han <fujun.han@iluvatar.ai>
2021-03-04 11:05:16 +08:00
0f1056390d Create PUBLICATIONS.md (#189) 2021-03-03 11:17:40 -08:00
34a42e5620 Update generator.py (#192) 2021-03-02 12:21:48 -08:00
8f09b82b12 Merge pull request #187 from NVIDIA/cutlass_2.5
CUTLASS 2.5.0
2021-02-26 23:56:04 -06:00
200a5a5146 Enabled reduction unit tests. 2021-02-26 15:46:57 -05:00
746b7b3247 Enabled tensor reduction kernels. 2021-02-26 15:32:19 -05:00
abdf16a4d9 Updated release notes. 2021-02-26 13:55:04 -05:00
0e13748649 CUTLASS 2.5 2021-02-26 09:58:26 -05:00
ccb697bac7 cutlass 2.4 documentation only update 2020-11-23 06:59:45 -06:00
e6bcdc60cf fix broken links (#148) 2020-11-19 21:46:54 -08:00
6615010cd0 CUTLASS 2.4 (Implicit GEMM convolution) (#147)
CUTLASS 2.4 (Implicit GEMM Convolution)

Co-authored-by: Manish Gupta <manigupta@nvidia.com>, Haicheng Wu <haichengw@nvidia.com>, Dustyn Blasig <dblasig@nvidia.com>, Andrew Kerr <akerr@nvidia.com>
2020-11-19 21:25:25 -08:00
c2b80ad4e4 Merge pull request #135 from NVIDIA/cutlass_2.3_final
CUTLASS 2.3.0
2020-09-25 13:25:26 -05:00
37a8f9e598 CUTLASS 2.3.0 final. 2020-09-25 10:34:46 -07:00
c53f3339bb CUTLASS 2.3 initial commit (#134)
CUTLASS 2.3 adds GEMMs targeting Sparse Tensor Cores on the NVIDIA Ampere Architecture, fast SGEMM, and small matrix classes, bug fixes, and performance enhancements.
2020-09-23 14:00:58 -07:00
4dac7490e6 Typoes (#107)
* Update splitk_gemm.cu

* Update gemm_bias_relu.cu

* Update mma_sm75.h
2020-07-13 14:25:52 -07:00
fd7e058d0c Added examples to enable the unity build (#102)
* Updated documentation of fused GEMM example and removed UNITY BUILD batch size. The default batch size when unity build is enabled tends to be favorable.
2020-06-17 07:09:18 -07:00
1ab1027954 Updated mma_sm80.h to avoid perf penalty due to reinterpret_cast<>. (#100)
- Updated mma_sm80.h to avoid perf penalty due to reinterpret_cast<>.
- Enhancement to CUTLASS Utility Library's HostTensorPlanarComplex template to support copy-in and copy-out
- Added test_examples target to build and test all CUTLASS examples
- Minor edits to documentation to point to GTC 2020 webinar
2020-06-15 10:47:01 -07:00
86931fef85 CUTLASS 2.2 (#96)
Adds support for NVIDIA Ampere Architecture features. CUDA 11 Toolkit recommended.
2020-06-08 16:17:35 -07:00
e33d90b361 update tools/library/CMakeLists to require python 3.6 according to #70 (#82)
#70 only updates the documentation. This commit reflects this bump in python version to the CMake configuration as well.
2020-04-08 10:54:36 -07:00
96dab34ad9 CUTLASS 2.1 (#83)
CUTLASS 2.1 contributes:
- BLAS-style host-side API added to CUTLASS Library
- Planar Complex GEMM kernels targeting Volta and Turing Tensor Cores
- Minor enhancements and bug fixes
2020-04-07 13:51:25 -07:00
7c0cd26d13 Need Python 3.6 to use enum.auto() (#70) 2019-11-22 09:39:12 -08:00
45ecbc885b Removed redundant conjugation operations from matrix_traits. (#65) 2019-11-20 11:27:13 -08:00
8aca98f9a7 Improved formatting, clarity, and content of several documents. (#64)
* Improved formatting, clarity, and content of several documents.
2019-11-20 10:42:15 -08:00
f4d9c8f755 Clang GPU compilation requires explicit CUDACC version flags (#63) 2019-11-20 09:52:11 -08:00
fb335f6a5f CUTLASS 2.0 (#62)
CUTLASS 2.0

Substantially refactored for

- Better performance, particularly for native Turing Tensor Cores
- Robust and durable templates spanning the design space
- Encapsulated functionality embodying modern C++11 programming techniques
- Optimized containers and data types for efficient, generic, portable device code

Updates to:
- Quick start guide
- Documentation
- Utilities
- CUTLASS Profiler

Native Turing Tensor Cores
- Efficient GEMM kernels targeting Turing Tensor Cores
- Mixed-precision floating point, 8-bit integer, 4-bit integer, and binarized operands

Coverage of existing CUTLASS functionality:
- GEMM kernels targeting CUDA and Tensor Cores in NVIDIA GPUs
- Volta Tensor Cores through native mma.sync and through WMMA API
- Optimizations such as parallel reductions, threadblock rasterization, and intra-threadblock reductions
- Batched GEMM operations
- Complex-valued GEMMs

Note: this commit and all that follow require a host compiler supporting C++11 or greater.
2019-11-19 16:55:34 -08:00
b5cab177a9 Performance enhancement for Volta Tensor Cores TN layout (#53)
* Fixed performance defect with indirect access to pointer array for Volta TensorCores TN arrangement.

* Updated patch version and changelog.

* Updated patch version and changelog.

* Added link to changelog in readme.

* Fixed markdown link
2019-07-10 10:54:12 -07:00
eb41735933 Merge pull request #47 from Artem-B/cutlass-1.3-clang
Make CUTLASS compileable with Clang.
2019-05-13 10:52:45 -07:00
fb8b3a98b7 Addressed code review comments. 2019-05-10 10:24:52 -07:00
d9d357877f Added missing file (#48) 2019-05-09 14:07:52 -07:00
e18292db46 Make CUTLASS compileable with Clang.
Requires a recent clang build (r359248 or newer).

Enable compilation with clang with these options:
cmake -DCUDA_COMPILER=clang -DCMAKE_CXX_COMPILER=/path/to/clang++
2019-05-02 11:00:22 -07:00
fe3438a3c1 cutlass 1.3.1 (#46)
CUTLASS 1.3.1 patch resolves failing text with NVRTC.
2019-04-19 16:54:52 -07:00
877bdcace6 Cutlass 1.3 Release (#42)
CUTLASS 1.3 Release
- Efficient GEMM kernel targeting Volta Tensor Cores via mma.sync instruction added in CUDA 10.1.
2019-03-20 10:49:17 -07:00
19a9d64e3c Removed patch version from README.
Removed patch version from README.
2018-12-19 15:20:43 -08:00
80e6f7c860 Merge pull request #38 from NVIDIA/resolve_maxwell
Resolved issue for incorrect SGEMM on Maxwell architecture.
2018-12-19 15:17:41 -08:00
822b0952cd Resolved issue for incorrect SGEMM on Maxwell architecture. 2018-12-19 15:07:16 -08:00
ed2ed4d667 Merge pull request #33 from NVIDIA/cutlass_1.2
CUTLASS 1.2
2018-10-26 14:59:50 -07:00
4db423c40f Minor edit to CHANGELOG. 2018-10-26 14:58:31 -07:00
b2bc0d3b79 Updating Doxygen docs 2018-10-26 14:54:58 -07:00
74df0331f2 CUTLASS 1.2 2018-10-26 14:38:46 -07:00
2332df492e Merge pull request #30 from NVIDIA/fix_utilities_example
Fixed cutlass_utilities example.
2018-09-29 15:09:18 -07:00
cfe4b933ef CUDA 9 lacks host-side conversions from float=>half. Instead, we must reinterpret_cast<> from cutlass::half_t => half. 2018-09-29 15:04:20 -07:00
6877595a5e Merge pull request #28 from NVIDIA/cutlass_1.1
Fixed typeo
2018-09-28 12:59:49 -07:00
69e3709da4 Fixed typeo
Fixed typeo
2018-09-28 12:59:20 -07:00
d419094c28 Merge pull request #26 from NVIDIA/cutlass_1.1
Clarification to README
2018-09-21 11:44:47 -07:00
1a7ac522f8 Clarification to README 2018-09-20 11:04:03 -07:00
bf6eec53eb Merge pull request #25 from NVIDIA/cutlass_1.1
Updated CUTLASS.md
2018-09-19 21:33:04 -07:00
206e38dac5 Updated copyright of CUTLASS.md 2018-09-19 21:31:12 -07:00
d85f6a1cec Merge pull request #24 from NVIDIA/cutlass_1.1
Cutlass 1.1
2018-09-19 21:16:53 -07:00
0826572c4c Reduced range of random values to avoid bit-level inconsistencies for large matrices. 2018-09-19 21:11:48 -07:00
77d1e0ca81 Updated README and CHANGELOG. 2018-09-19 20:42:51 -07:00
d7137f9c0a Updated doxygen 2018-09-19 14:02:08 -07:00
461f417b9d Checkpointing CUTLASS 1.1 release. 2018-09-18 16:58:03 -07:00
cf0301e00f Merge pull request #15 from NVIDIA/release_1.0.1_edits
Minor edits to README and changelog pursuant CUTLASS 1.0.1 patch.
2018-06-26 13:59:01 -07:00
b9bb0d1a49 Edits to README and changelog pursuant CUTLASS 1.0.1 patch. 2018-06-26 13:57:39 -07:00
e1c4ba501b Merge pull request #13 from NVIDIA/cutlass_v1.0.1
Cutlass v1.0.1
2018-06-12 08:25:56 -07:00
c566e83e6d Updated changelog. 2018-06-11 14:54:07 -07:00
374882be53 Replaced GoogleTest copy with submodule. Added updates to support intra-threadblock reductions. Added tests for same. 2018-06-11 11:47:15 -07:00
2c496c3e9e Replaced GoogleTest copy with Git submodule. 2018-06-11 11:32:41 -07:00
9fd55460c6 Merge pull request #10 from NVIDIA/cutlass_v1.0_rel
Minor updates to usage and README.
2018-05-18 12:27:31 -07:00
480732c2e8 Minor updates to usage and readme. 2018-05-17 15:10:55 -07:00
68aaee8773 Merge pull request #9 from NVIDIA/cutlass_v1.0_rel
Updated URL to Doxygen and modified usage statement
2018-05-17 11:12:37 -07:00
acb90e962a Updated url to Doxygen and modified usage statement in performance test program. 2018-05-17 11:11:05 -07:00
96bc3f227f Merge pull request #8 from NVIDIA/cutlass_v1.0_rel
Configured Github Pages
2018-05-16 15:26:55 -07:00
25ff282403 Moved Doxygen documents. 2018-05-16 15:25:24 -07:00
9d5726a568 Set theme jekyll-theme-minimal 2018-05-16 13:49:06 -07:00
6f0d271d8d CUTLASS v1.0
CUTLASS v1.0 released.
2018-05-16 13:47:13 -07:00
923dfb42ce Updated README.md 2018-05-16 12:50:10 -07:00
6f6f269a0a Updated README.md 2018-05-16 12:47:07 -07:00
2028ebe120 CUTLASS v1.0 release 2018-05-16 11:44:56 -07:00
84377249a1 Merge pull request #2 from Artem-B/clang-fixes
Merging "Clang fixes" into master.
2018-01-04 15:52:53 -08:00
901287175f Merge branch 'Artem-B-clang-fixes' 2018-01-04 15:46:08 -08:00
1c9b54df16 Whitespace fix. 2018-01-03 16:42:51 -08:00
39616514d0 Reworked CUDA_LOG macro to print location&the message with one printf.
This replies on the fact that clang allows using device-side features
from __host__/__device__ functions from __host__ ones as long as we
don't have to generate code for that. Wrapping thread/blockIdx in
__host__ __device__ function allows using CUDA_LOG everywhere during
host and device compilation.
2018-01-03 16:36:50 -08:00
df4b4e4bb6 Added _cuda_ to the name of the executable to indicate that it's not clang's version. 2017-12-11 16:34:10 -08:00
81957b3a3d Force inlining of few functions that rely on that for performance.
Clang is less agressive than nvccnvcc, so number of functions did not getn
inlined into the kernel by default. That prevented SROA from eliminating
loads/stores to temporary buffers and resulted in abysmal performance.

Replaced inline with __forceinline__ to ensure that we do inline the
functions necessary for optimal performance.
2017-12-11 14:52:30 -08:00
ce2b3f695d Fixed debug macros for clang.
Unlike nvcc, clang always sees both host and device-side code during
compilation. CUDA_LOG macro is used in both host and device code, so when it
expanded to contain device-only code, that resulted in errors when it was used
from the host-side functions.

In order to make CUDA_LOG work with clang it was split into two parts -- a pair
of target-attribute-based overloaded functions that perform host or device
specific parts of logging, and a printf which works on both sides.
2017-12-11 14:52:30 -08:00
e9e7cd4d44 Make cutlass compilable with clang.
E.g:
PATH=/nvcc/path/bin:/clang/path/bin:$PATH make sm=35,60 compiler=clang all
2017-12-11 14:52:30 -08:00
95b0578d34 Update license info 2017-12-06 10:00:59 -05:00
f4b48c7669 Update README.md 2017-12-05 22:58:46 -05:00
6cb88d53eb Update README.md 2017-12-05 22:58:12 -05:00
537a4bcedf Update README.md 2017-12-05 22:54:49 -05:00
5bd3f09312 Update README.md 2017-12-05 22:53:11 -05:00
6f091f5620 Update README.md 2017-12-05 22:44:01 -05:00
0428c89fd5 Updating readme with relative per chart 2017-12-05 22:40:47 -05:00
e2bf51c3fe Update README.md 2017-12-05 22:25:42 -05:00
57747e382e Update README.md 2017-12-05 21:32:06 -05:00
dd4dd4cebf Update README.md 2017-12-05 20:58:01 -05:00
6565b48747 Update README.md 2017-12-05 20:56:49 -05:00
73211bbb88 Update README.md 2017-12-05 20:55:54 -05:00
9dcb2b4c7d Update README.md 2017-12-05 20:55:03 -05:00
f30abfc00a Update README.md 2017-12-05 20:50:15 -05:00
8ebd6b06d0 Replace svg with png+text 2017-12-05 20:20:25 -05:00
04ffa156e8 Adding figure to readme.md 2017-12-05 20:15:33 -05:00
24d0ba65c5 Update code formatting 2017-12-05 15:51:01 -05:00
5369 changed files with 1133420 additions and 10963 deletions

23
.github/ISSUE_TEMPLATE/bug_report.md vendored Normal file
View File

@ -0,0 +1,23 @@
---
name: Bug report
about: Create a bug report to help us improve CUTLASS
title: "[BUG]"
labels: "? - Needs Triage, bug"
assignees: ''
---
**Describe the bug**
A clear and concise description of what the bug is.
**Steps/Code to reproduce bug**
Follow this guide http://matthewrocklin.com/blog/work/2018/02/28/minimal-bug-reports to craft a minimal bug report. This helps us reproduce the issue you're having and resolve the issue more quickly.
**Expected behavior**
A clear and concise description of what you expected to happen.
**Environment details (please complete the following information):**
- Environment location: [Bare-metal, Docker, Cloud(specify cloud provider)]
**Additional context**
Add any other context about the problem here.

View File

@ -0,0 +1,35 @@
---
name: Documentation request
about: Report incorrect or needed documentation to improve CUTLASS
title: "[DOC]"
labels: "? - Needs Triage, documentation"
assignees: ''
---
## Report incorrect documentation
**Location of incorrect documentation**
Provide links and line numbers if applicable.
**Describe the problems or issues found in the documentation**
A clear and concise description of what you found to be incorrect.
**Steps taken to verify documentation is incorrect**
List any steps you have taken:
**Suggested fix for documentation**
Detail proposed changes to fix the documentation if you have any.
---
## Report needed documentation
**Report needed documentation**
A clear and concise description of what documentation you believe it is needed and why.
**Describe the documentation you'd like**
A clear and concise description of what you want to happen.
**Steps taken to search for needed documentation**
List any steps you have taken:

View File

@ -0,0 +1,20 @@
---
name: Feature request
about: Suggest an idea for CUTLASS
title: "[FEA]"
labels: "? - Needs Triage, feature request"
assignees: ''
---
**Is your feature request related to a problem? Please describe.**
A clear and concise description of what the problem is. Ex. I wish I could use CUTLASS to do [...]
**Describe the solution you'd like**
A clear and concise description of what you want to happen.
**Describe alternatives you've considered**
A clear and concise description of any alternative solutions or features you've considered.
**Additional context**
Add any other context, code examples, or references to existing implementations about the feature request here.

View File

@ -0,0 +1,10 @@
---
name: Submit question
about: Ask a general question about CUTLASS
title: "[QST]"
labels: "? - Needs Triage, question"
assignees: ''
---
**What is your question?**

11
.github/workflows/labeler.yml vendored Normal file
View File

@ -0,0 +1,11 @@
name: "Pull Request Labeler"
on:
- pull_request_target
jobs:
triage:
runs-on: ubuntu-latest
steps:
- uses: actions/labeler@main
with:
repo-token: "${{ secrets.GITHUB_TOKEN }}"

View File

@ -0,0 +1,35 @@
name: Auto Assign New Issues to Triage Project
on:
issues:
types: [opened]
env:
GITHUB_TOKEN: ${{ secrets.GITHUB_TOKEN }}
jobs:
assign_one_project:
runs-on: ubuntu-latest
name: Assign to New Issues to Triage Project
steps:
- name: Process bug issues
uses: docker://takanabe/github-actions-automate-projects:v0.0.1
if: contains(github.event.issue.labels.*.name, 'bug') && contains(github.event.issue.labels.*.name, '? - Needs Triage')
env:
GITHUB_TOKEN: ${{ secrets.GITHUB_TOKEN }}
GITHUB_PROJECT_URL: https://github.com/NVIDIA/cutlass
GITHUB_PROJECT_COLUMN_NAME: 'Needs prioritizing'
- name: Process feature issues
uses: docker://takanabe/github-actions-automate-projects:v0.0.1
if: contains(github.event.issue.labels.*.name, 'feature request') && contains(github.event.issue.labels.*.name, '? - Needs Triage')
env:
GITHUB_TOKEN: ${{ secrets.GITHUB_TOKEN }}
GITHUB_PROJECT_URL: https://github.com/NVIDIA/cutlass
GITHUB_PROJECT_COLUMN_NAME: 'Needs prioritizing'
- name: Process other issues
uses: docker://takanabe/github-actions-automate-projects:v0.0.1
if: contains(github.event.issue.labels.*.name, '? - Needs Triage') && (!contains(github.event.issue.labels.*.name, 'bug') && !contains(github.event.issue.labels.*.name, 'feature request'))
env:
GITHUB_TOKEN: ${{ secrets.GITHUB_TOKEN }}
GITHUB_PROJECT_URL: https://github.com/NVIDIA/cutlass
GITHUB_PROJECT_COLUMN_NAME: 'Needs prioritizing'

57
.github/workflows/stale.yml vendored Normal file
View File

@ -0,0 +1,57 @@
name: Mark inactive issues and pull requests
on:
schedule:
- cron: "0 * * * *"
jobs:
mark-inactive-30d:
runs-on: ubuntu-latest
steps:
- name: Mark 30 day inactive issues and pull requests
uses: actions/stale@v3
with:
repo-token: ${{ secrets.GITHUB_TOKEN }}
stale-issue-message: >
This issue has been labeled `inactive-30d` due to no recent activity in the past 30 days.
Please close this issue if no further response or action is needed.
Otherwise, please respond with a comment indicating any updates or changes to the original issue and/or confirm this issue still needs to be addressed.
This issue will be labeled `inactive-90d` if there is no activity in the next 60 days.
stale-issue-label: "inactive-30d"
exempt-issue-labels: "0 - Blocked,0 - Backlog,good first issue"
days-before-issue-stale: 30
days-before-issue-close: -1
stale-pr-message: >
This PR has been labeled `inactive-30d` due to no recent activity in the past 30 days.
Please close this PR if it is no longer required.
Otherwise, please respond with a comment indicating any updates.
This PR will be labeled `inactive-90d` if there is no activity in the next 60 days.
stale-pr-label: "inactive-30d"
exempt-pr-labels: "0 - Blocked,0 - Backlog,good first issue"
days-before-pr-stale: 30
days-before-pr-close: -1
operations-per-run: 50
mark-inactive-90d:
runs-on: ubuntu-latest
steps:
- name: Mark 90 day inactive issues and pull requests
uses: actions/stale@v3
with:
repo-token: ${{ secrets.GITHUB_TOKEN }}
stale-issue-message: >
This issue has been labeled `inactive-90d` due to no recent activity in the past 90 days.
Please close this issue if no further response or action is needed.
Otherwise, please respond with a comment indicating any updates or changes to the original issue and/or confirm this issue still needs to be addressed.
stale-issue-label: "inactive-90d"
exempt-issue-labels: "0 - Blocked,0 - Backlog,good first issue"
days-before-issue-stale: 90
days-before-issue-close: -1
stale-pr-message: >
This PR has been labeled `inactive-90d` due to no recent activity in the past 90 days.
Please close this PR if it is no longer required.
Otherwise, please respond with a comment indicating any updates.
stale-pr-label: "inactive-90d"
exempt-pr-labels: "0 - Blocked,0 - Backlog,good first issue"
days-before-pr-stale: 90
days-before-pr-close: -1
operations-per-run: 50

2
.gitignore vendored Normal file
View File

@ -0,0 +1,2 @@
# PyCache files
__pycache__/

0
.gitmodules vendored Normal file
View File

330
CHANGELOG.md Normal file
View File

@ -0,0 +1,330 @@
# NVIDIA CUTLASS Changelog
## [3.0.0](https://github.com/NVIDIA/cutlass/releases/tag/v3.0.0) (2023-01-23)
* [CuTe](/media/docs/cute/00_quickstart.md), a [new core library and backend](/include/cute) for CUTLASS 3.0 that defines a single Layout vocabulary type and an associated algebra of layouts for a much more expressive and composable abstraction for tensors, sets of parallel agents, and operations by said agents on tensors.
* [A new conceptual operation hierarchy](media/docs/cutlass_3x_design.md) that replaces the architecture-centric hierarchy of CUTLASS 2.x and [documentation for CUTLASS 3.0's GEMM API changes](/media/docs/gemm_api_3x.md).
* Strict API backwards compatibility that exposes both 2.x and 3.x API kernels through the same [`device::GemmUniversalAdapter`](include/cutlass/gemm/device/gemm_universal_adapter.h) and [`kernel::GemmUniversal`](include/cutlass/gemm/kernel/gemm_universal.hpp) types, allowing users to include both APIs in the same translation units. More information can be found in the [3.x backwards compatibility section](media/docs/cutlass_3x_backwards_compatibility.md).
* Updates to [Functionality](media/docs/functionality.md) which directs users on which kernels are supported via CUTLASS-2 and CUTLASS-3.
* Updates to [Compatibility](/README.md#compatibility) Section regarding supported compilers, operating systems, CUDA Toolkits, Hardware Architectures and [Target Architecture](/README.md#Target-Architecture).
* New warp-specialized GEMM [kernel schedules](include/cutlass/gemm/kernel/sm90_gemm_tma_warpspecialized.hpp) and [mainloops](include/cutlass/gemm/collective/sm90_mma_tma_gmma_ss_warpspecialized.hpp) targeting Hopper architecture that achieve great performance with TMA, WGMMA, and threadblock clusters.
* Extensions to CUTLASS profiler to support threadblock cluster shapes in library and profiler tile configurations.
* [CUTLASS library integration](/tools/library/src/gemm_operation_3x.hpp) for 3.x API kernels built through the new `CollectiveBuilder` API, enabling CUTLASS profiler.
* Support for [Hopper GEMMs](examples/48_hopper_warp_specialized_gemm) through the new 3.0 API with CuTe-based exposure of the Hopper [Tensor Memory Accelerator](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-cp-async-bulk-tensor) and [WGMMA Tensor Core](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#asynchronous-warpgroup-level-matrix-instructions) features.
* Set of examples that demonstrate the usage of the new 3.0 API to easily build GEMM kernels targeting Hopper: examples [48](examples/48_hopper_warp_specialized_gemm), [49](examples/49_hopper_gemm_schedules_with_collective_builder), and [50](examples/50_hopper_gemm_with_epilogue_swizzle).
## [2.11.0](https://github.com/NVIDIA/cutlass/releases/tag/v2.11.0) (2022-11-19)
* [Stream-K](/examples/47_ampere_gemm_universal_streamk), which is a new general way to do split-K. It can not only improve performance, but can also significantly reduce the number of tile sizes that need to be profiled to find the best one.
* [Fused multi-head attention Kernel](/examples/41_fused_multi_head_attention). It has two variants: one uses batched GEMM for the fixed sequence length, and the other one uses group GEMM for the variable sequence length. Both versions just need one kernel.
* [Dual GEMM](/examples/45_dual_gemm), which can fuse A x B and A x C into one kernel. Two GEMMs has no producer-consumer dependency.
* Hopper improves [double precision matrix multiplication](/test/unit/gemm/device/gemm_f64n_f64t_f64t_tensor_op_f64_sm90.cu) by 2x compared to Ampere at iso-clocks. It is supported since CUDA 11.8.
* [BLAS3](/test/unit/gemm/device/hemm_cf64_cf64_cf64_tensor_op_f64_sm90.cu) functions with Hoppers new double precision matrix multiplication instructions.
* [ELL Block Sparse GEMM](/examples/43_ell_block_sparse_gemm), which uses an [ELL matrix](https://developer.nvidia.com/blog/accelerating-matrix-multiplication-with-block-sparse-format-and-nvidia-tensor-cores/) to describe the sparsity of A matrix. B and output matrices are still dense. The block size can be arbitary.
* Optimized [Group Conv](/examples/42_ampere_tensorop_group_conv) for SingleGroup mode, which requires that the output channel per group is a multiple of Threadblock tile N.
* [Optimized DepthWise Conv](/examples/46_depthwise_simt_conv2dfprop/depthwise_simt_conv2dfprop.cu). Two new modes are added
* [kOptimized](/test/unit/conv/device/depthwise_conv2d_fprop_direct_conv_f16nhwc_f16nhwc_f16nhwc_simt_f16_sm60.cu) - use direct conv to compute instead of implicit GEMM.
* The restrictions are: 1) input ,output channel and group number should be multiple of (128 / sizeof(input element)). 2) The input filter size should be the same as the template parameter configuration.
* [kFixedStrideDilation](/test/unit/conv/device/depthwise_conv2d_fprop_direct_conv_fixed_stride_dilation_f16nhwc_f16nhwc_f16nhwc_simt_f16_sm60.cu) - which puts stride and dilation into templates to further improve the performance. In this mode, kernel persistents some inputs into register to squeeze more performance, so large filter/stride/dilation is not recommanded.
* The restrictions are: 1) input, output channel and group number should be multiple of (128 / sizeof(input element)). 2) input filter size, stride, dilation should same as the template parameter configuration.
* [Scripts](/examples/44_multi_gemm_ir_and_codegen) to fuse multiple back-to-back GEMM. Its implementation was discussed in a GTC'22 Spring [talk](https://www.nvidia.com/en-us/on-demand/session/gtcspring22-s41606/).
* [FP8 data type definition](/include/cutlass/float8.h) and [conversion routines](/include/cutlass/numeric_conversion.h#L1274-2115).
* Updates and bugfixes from the community (thanks!). Big shout out to Meta's [xFormers](https://github.com/facebookresearch/xformers).
* **Deprecation announcement:** CUTLASS plans to deprecate the following:
* Maxwell and Pascal GPU architectures
* Ubuntu 16.04
* CUDA 10.2
## [2.10.0](https://github.com/NVIDIA/cutlass/releases/tag/v2.10.0) (2022-08-23)
* [CUTLASS Python](/examples/40_cutlass_py) now supports GEMM, CONV, Group GEMM for different data types as well as different epilogue flavours.
* Optimizations for CUTLASS's [Grouped GEMM](examples/24_gemm_grouped/gemm_grouped.cu) kernel. Threadblock scheduling part is improved. Some computation can be moved to the host side if applicable. [Grouped Syr2k](examples/38_syr2k_grouped/syr2k_grouped.cu) kernels are added, too.
* Optimizations for [GEMM+Softmax](examples/35_gemm_softmax). All the reduction computation is fused into the previous GEMM. More template arguments are provided to fine tune the performance.
* [Grouped GEMM for Multihead Attention](examples/41_multi_head_attention). This general group gemm based MHA does not require the sequence length of all GEMMs to be the same which makes it most useful for natural language processing.
* [GEMM + Layer norm fusion for Ampere](examples/37_gemm_layernorm_gemm_fusion/) splits the layernorm into two parts and both of them can be fused into the GEMMs before and after separately. In addition to use square sum to compute variance of layernorm, [Shift-K](https://en.wikipedia.org/wiki/Algorithms_for_calculating_variance#Computing_shifted_data) is provided if square sum raise numerical issues.
* [GEMM Epilogue Permutation Fusion](examples/39_gemm_permute) can apply user provided permutation layout mapping in the GEMM epilogue.
* [Grouped convolution targeting implicit GEMM](test/unit/conv/device/group_conv2d_fprop_implicit_gemm_f16nhwc_f16nhwc_f16nhwc_tensor_op_f32_sm80.cu) introduces the first group convolution implementation to CUTLASS. It is an Analytical implementation, not an Optimized. The restrictions are: 1) input and output channel number should be multiple of group number. 2) split-K is not supported. The implementation has 2 modes:
* kSingleGroup: output channel per group is multiple of Threadblock tile N.
* kMultipleGroup: Threadblock tile N is multiple of output channel per group.
* [Depthwise separable convolution](test/unit/conv/device/depthwise_fprop_implicit_gemm_f16nhwc_f16nhwc_f16nhwc_simt_f16_sm60.cu) introduces the first depthwise convolution which is also Analytical for now. The restrictions are: 1) SIMT only 2) No split-K 3) input channel equals to output channel equals to group number.
* Standalone [Layernorm](/tools/util/include/cutlass/util/device_layernorm.h) and [Pooling](/tools/util/include/cutlass/util/device_nhwc_pooling.h) kernels.
* [Back-to-back GEMM/CONV](examples/13_two_tensor_op_fusion) relaxes the requirement that the first GEMM K dimension needs to be the multiple of Threadblock Tile K dimension.
* Optimal performance using [**CUDA 11.6u2**](https://developer.nvidia.com/cuda-downloads)
* Updates and bugfixes from the community (thanks!)
## [2.9.0](https://github.com/NVIDIA/cutlass/releases/tag/v2.9.0) (2022-04-21)
* [First layer Convolution kernels](/test/unit/conv/device/conv2d_fprop_fixed_channels_f16nhwc_f16nhwc_f16nhwc_tensor_op_f32_sm80.cu) specialized for small channel counts and reduced alignment
* [Few channels](/include/cutlass/conv/threadblock/conv2d_fprop_activation_tile_access_iterator_few_channels.h) specialization for reduced alignment capabilities
* [Fixed channels](/include/cutlass/conv/threadblock/conv2d_fprop_activation_tile_access_iterator_fixed_channels.h) further specialized when channel count perfectly matches the access vector size
* [Unit tests](/test/unit/conv/device/conv2d_fprop_few_channels_f16nhwc_f16nhwc_f16nhwc_tensor_op_f32_sm80.cu)
* [Python-based instance emitter](/tools/library/scripts/generator.py) in the CUTLASS Library and support in the Profiler
* [BLAS3](https://docs.nvidia.com/cuda/cublas/index.html#cublas-level-3-function-reference) operators accelerated by Tensor Cores
* Supported types: f32, cf32, f64, cf64, tf32x3, complex tf32x3
* [HERK](/test/unit/gemm/device/her2k_cf32h_cf32n_tensor_op_fast_f32_sm80.cu) with [emitter](/tools/library/scripts/rank_k_operation.py)
* [SYRK](/test/unit/gemm/device/syrk_f32n_f32t_tensor_op_fast_f32_sm80.cu) with [emitter](/tools/library/scripts/rank_k_operation.py)
* [SYMM](/test/unit/gemm/device/symm_f32n_f32n_tensor_op_fast_f32_ls_sm80.cu) with [emitter](/tools/library/scripts/symm_operation.py)
* [TRMM](/test/unit/gemm/device/trmm_f32n_f32t_f32t_tensor_op_fast_f32_ls_sm80.cu) with [emitter](/tools/library/scripts/trmm_operation.py)
* [Unit tests](/test/unit/gemm/device/testbed_rank_k_universal.h)
* [CUTLASS Python](/examples/40_cutlass_py) demonstrating JIT compilation of CUTLASS kernels and a Python-based runtime using [CUDA Python](https://developer.nvidia.com/cuda-python)
* [Python-based runtime](/tools/library/scripts/rt.py) interoperable with existing emitters
* [GEMM + Softmax example](/examples/35_gemm_softmax)
* [Gather and Scatter Fusion with GEMM](/examples/36_gather_scatter_fusion) can gather inputs and scatters outputs based on indices vectors in the same GEMM kernel.
* It can select random rows in a row major matrix.
* It can select random columns in a column major matrix.
* [Back-to-back GEMM/CONV](examples/13_two_tensor_op_fusion) fully supports buffering the first GEMM/CONV results in the shared memory for the latter one to use. It can eliminate register spill when the tile size is big. Additionally, bias vector add is supported in the first GEMM/CONV.
* Supported kernels: GEMM and CONV.
* Supported types: fp16 and int8.
* Supported architectures: Turing and Ampere.
* [Transposed Convolution](/examples/34_transposed_conv2d) (a.k.a Deconvolution) support which reuses Dgrad implementation.
* [Utility functions](/tools/util/include/cutlass/util) that can pad NHWC and convert between NCHW and NHWC.
* [Small alignment implicit gemm](https://github.com/NVIDIA/cutlass/issues/242) support for Fprop/Dgrad/Wgrad so that padding is no longer mandated to use tensor cores in these kernels.
* Epilogue enhancement:
* Eliminate bank conflicts in int8 tensor core kernels.
* Half2 usage if epilogue compute type is fp16.
* More activation functions: Silu, Hardswish, Leaky Relu.
* New elementwise fusion pattern for [residual block](/include/cutlass/epilogue/thread/linear_combination_residual_block.h).
* [Group GEMM](/examples/24_gemm_grouped) thread block number calculation fix which helps to launch the intended number of threadblocks to fully occupy the GPUs.
* [Parallel GEMM splitk](https://github.com/NVIDIA/cutlass/pull/277) support in the CUTLASS profiler.
* Optimal performance using [**CUDA 11.6u2**](https://developer.nvidia.com/cuda-downloads)
* Updates and bugfixes from the community (thanks!)
## [2.8.0](https://github.com/NVIDIA/cutlass/releases/tag/v2.8.0) (2021-11-19)
* **TF32x3:** emulated single-precision using Tensor Cores
* 45+ TFLOPs on NVIDIA A100
* [GEMM SDK example](/examples/27_ampere_3xtf32_fast_accurate_tensorop_gemm/27_ampere_3xtf32_fast_accurate_tensorop_gemm.cu) (real)
* [COMPLEX GEMM SDK example](/examples/29_ampere_3xtf32_fast_accurate_tensorop_complex_gemm/29_ampere_3xtf32_fast_accurate_tensorop_complex_gemm.cu) (complex)
* [Implicit GEMM Convolution SDK example](/examples/28_ampere_3xtf32_fast_accurate_tensorop_fprop/ampere_3xtf32_fast_accurate_tensorop_fprop.cu)
* **Mainloop fusion for Convolution:** convolution with fused per-channel scale-bias-relu
* [Conv Fprop SDK example](/examples/25_ampere_fprop_mainloop_fusion/ampere_fprop_mainloop_fusion.cu)
* [Conv WGrad SDK example](/examples/26_ampere_wgrad_mainloop_fusion/ampere_wgrad_mainloop_fusion.cu)
* [cutlass::conv::device::ImplicitGemmConvolutionFusion](/include/cutlass/conv/device/implicit_gemm_convolution_fusion.h)
* **Grouped GEMM:** similar to batched GEMM with distinct problem size per group
* [SDK example](/examples/24_gemm_grouped) with performance comparison with Batched Strided GEMM
* [cutlass::gemm::device::GemmGrouped](/include/cutlass/gemm/device/gemm_grouped.h)
* [Implicit GEMM Convolution fusion](/examples/13_two_tensor_op_fusion/) supports staging 1st convolution's output accumulator in the shared memory on Turing. This allows more flexible warp tile sizes and less regsiter pressue.
* Optimal performance using [**CUDA 11.5**](https://developer.nvidia.com/cuda-downloads)
* Updates from the community (thanks!)
* **Deprecation announcement:** CUTLASS plans to deprecate the following:
* Maxwell and Pascal GPU architectures
* Ubuntu 16.04
* CUDA 10.2
## [2.7.0](https://github.com/NVIDIA/cutlass/releases/tag/v2.7.0) (2021-09-24)
* Mainloop fusion for GEMM: [summation over A or B](/examples/23_ampere_gemm_operand_reduction_fusion/ampere_gemm_operand_reduction_fusion.cu)
* [Strided DGRAD (optimized iterators)](/include/cutlass/conv/kernel/default_conv2d_dgrad.h)
* [Half-precision GELU_taylor activation functions](/include/cutlass/epilogue/thread/activation.h#L196)
* Use these when accumulation and epilogue compute types are all `cutlass::half_t`
* Tuning and bug fixes to [fused GEMM + GEMM example](/examples/13_two_tensor_op_fusion/)
* Support for smaller than 128b aligned Convolutions: [see examples](test/unit/conv/device/conv2d_fprop_implicit_gemm_f16nhwc_f16nhwc_f16nhwc_tensor_op_f16_sm80.cu#L272)
* Caching of results to accelerate Convolution [unit tests](test/unit/conv/device/cache_testbed_output.h)
* Can be enabled or disabled by running `cmake .. -DCUTLASS_TEST_ENABLE_CACHED_RESULTS=OFF`
* Corrections and bug fixes reported by the CUTLASS community
* Thank you for filing these issues!
## [2.6.1](https://github.com/NVIDIA/cutlass/releases/tag/v2.6.1) (2021-09-03)
* Arbitrary padding and striding for CUTLASS Strided DGRAD Convolution operator (Analytic Iterators)
* Tuning for GEMMs fused with partial reductions
* Corrections and bug fixes reported by the CUTLASS community
* Thank you for filing these issues!
## [2.6.0](https://github.com/NVIDIA/cutlass/releases/tag/v2.6.0) (2021-07-22)
* Optimal performance when compiled with the [CUDA 11.4 Toolkit](https://developer.nvidia.com/cuda-toolkit)
* Adopt the new L2 prefetch feature in [cp.async](/include/cutlass/arch/memory.h) and [global load](/include/cutlass/arch/memory_sm80.h)
* Fused operators with GEMM and Convolution
* [Fused broadcast in epilogue](test/unit/gemm/device/gemm_with_broadcast_f16n_f16n_f16n_tensorop_f32_sm75.cu)
* [Fused partial reduction in epilogue](/test/unit/gemm/device/gemm_with_reduction_f16n_f16n_f16n_tensorop_f32_sm75.cu)
* 64b tensor strides and leading dimensions support for GEMMs
* Affine rank=2 matrix layouts
* Row stride and column stride for matrices using [cutlass::layout::AffineRank2](/include/cutlass/layout/matrix.h)
* Support [FP64 tensor core](/examples/18_ampere_fp64_tensorop_affine2_gemm/ampere_fp64_tensorop_affine2_gemm.cu) and SIMT GEMM.
* [Batched GEMV](/test/unit/gemm/device/gemv.cu) preview implementation
* [New strided Dgrad](test/unit/conv/device/conv2d_strided_dgrad_implicit_gemm_f16nhwc_f16nhwc_f32nhwc_tensor_op_f32_sm80.cu) implementation
* Accelerates over previous implementation by cutting down redundant math by 4x
* Support using new `Dy` and `w` analytic iterators and existing `cutlass::conv::device::ImplicitGemmConvolution` interface
* Quaternion-valued GEMM and Convolution in single- and double-precision (targeting CUDA Cores)
* Updates to [quaternion.h](/include/cutlass/quaternion.h) and [functional.h](/include/cutlass/functional.h)
* SDK Example for [GEMM](/examples/21_quaternion_gemm/quaternion_gemm.cu) and [Convolution](/examples/22_quaternion_gemm/quaternion_conv.cu)
* [Unit tests for GEMM](/test/unit/gemm/device/simt_qgemm_nn_sm50.cu) and [Convolution](/test/unit/conv/device/conv2d_fprop_implicit_gemm_qf32nhwc_qf32nhwc_qf32nhwc_simt_f32_sm50.cu)
* Many improvements to the epilogue.
* Provide an [option](/include/cutlass/epilogue/threadblock/epilogue.h) to not fully unroll the epilogue to reduce the code size and improve the performance when using complicated elementwise operations
* Performance improvement for FP16 tensor core kernels
* Bug fixes
* Enhanced Clang support and the combination of Clang 13 and CUDA 11.4 can build and run kernels from Pascal and Ampere.
* Updated minimum CUDA Toolkit requirement to 10.2
* [CUDA 11.4 Toolkit](https://developer.nvidia.com/cuda-toolkit) recommended
* Corrections and bug fixes reported by the CUTLASS community
* Thank you for filing these issues!
## [2.5.0](https://github.com/NVIDIA/cutlass/releases/tag/v2.5.0) (2021-02-26)
* Tensor reductions
* _m_-to-_n_ reductions of tensors with affine layout
* [Specializations](/test/unit/reduction/device/tensor_reduce_contiguous.cu) for reductions including contiguous dimension
* [Specializations](/test/unit/reduction/device/tensor_reduce_strided.cu) for reductions excluding contiguous dimension
* Custom reduction functors such as `cutlass::logical_and`
* Large tensor support, up to 2^63 elements (however, each dimension is limited to an extent of 2^31)
* Optimizations for 3-D convolution
* [Optimized tile iterators](include/cutlass/conv/threadblock/conv3d_fprop_activation_tile_access_iterator_optimized.h) using precomputed delta table for 3-D convolution
* Full coverage of [forward](test/unit/conv/device/conv3d_fprop_implicit_gemm_f16ndhwc_f16ndhwc_f32ndhwc_tensor_op_f32_sm80.cu) and [backwards](test/unit/conv/device/conv3d_dgrad_implicit_gemm_f16ndhwc_f16ndhwc_f32ndhwc_tensor_op_f32_sm80.cu) passes for 3D convolution
* [Fused Convolution+Convolution example](/examples/13_two_tensor_op_fusion/README.md)
* Corrections and bug fixes reported by the CUTLASS community
* Thank you for filing these issues!
## [2.4.0](https://github.com/NVIDIA/cutlass/releases/tag/v2.4.0) (2020-11-19)
* Implicit GEMM convolution kernels supporting CUDA and Tensor Cores on NVIDIA GPUs
* Operators: forward (Fprop), backward data gradient (Dgrad), and backward weight gradient (Wgrad) convolution
* Data type: FP32, complex<FP32>, Tensor Float 32 (TF32), BFloat16 (BF16), Float16, Int4, Int8, Int32
* Spatial dimensions: 1-D, 2-D, and 3-D
* Layout: NHWC, NCxHWx
* Implicit GEMM convolution components:
* Global memory iterators supporting Fprop, Dgrad, and Wgrad
* `MmaMultistage` for implicit GEMM convolution for NVIDIA Ampere architecture
* `MmaPipeline` for implicit GEMM convolution for NVIDIA Volta and Turing architectures
* [Documentation](/media/docs/implicit_gemm_convolution.md) describing Implicit GEMM Convolution algorithm and implementation
## [2.3.0](https://github.com/NVIDIA/cutlass/releases/tag/v2.3.0) (2020-09-23)
* [NVIDIA Ampere Architecture features](https://devblogs.nvidia.com/nvidia-ampere-architecture-in-depth/)
* [Sparse Tensor Core GEMM kernels](test/unit/gemm/device/gemm_f16n_f16n_f32t_tensor_op_f32_sparse_sm80.cu):
* Direct access to Sparse Tensor Cores and maximum performance via [`mma.sp.sync`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#warp-level-matrix-instructions-mma-and-friends)
* Fast SGEMM targeting GeForce RTX 30-series CUDA Cores
* Minor Features:
* [Activation functions](/include/cutlass/epilogue/thread/activation.h) such as [GeLU](/include/cutlass/epilogue/thread/linear_combination_gelu.h) and [Sigmoid](/include/cutlass/epilogue/thread/linear_combination_sigmoid.h)
* Small [matrix](/include/cutlass/matrix.h) and [quaternion](/include/cutlass/quaternion.h) template classes in device code
* [Floating-point constants](/include/cutlass/constants.h)
* NVIDIA Ampere GPU Architecture examples and documentation:
* [Tensor Float 32](/examples/14_ampere_tf32_tensorop_gemm/ampere_tf32_tensorop_gemm.cu) and
* [Sparse Tensor Cores](/examples/15_ampere_sparse_tensorop_gemm/ampere_sparse_tensorop_gemm.cu)
* Documentation added on CUTLASS [efficient row-major epilogue](/media/docs/gemm_api.md#efficient-epilogue)
## [2.2.0](https://github.com/NVIDIA/cutlass/releases/tag/v2.2.0) (2020-06-08)
* [NVIDIA Ampere Architecture features](https://devblogs.nvidia.com/nvidia-ampere-architecture-in-depth/)
* Fast Tensor Core operations:
* Maximum performance via [`mma.sync`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#warp-level-matrix-instructions-mma-and-friends)
* Tensor Float 32, BFloat16, and double-precision data types
* Mixed integer data types (int8, int4, bin1)
* Asynchronous copy for deep software pipelines via [`cp.async`](https://docs.nvidia.com/cuda/parallel-thread-execution)
* Described in [GTC 2020 Webinar (SR 21745)](https://developer.nvidia.com/gtc/2020/video/s21745) (free registration required)
* Features:
* SDK examples showing GEMM fused with bias+relu and fused GEMM+GEMM
* Complex-valued GEMMs targeting NVIDIA Ampere Tensor Cores in double-precision and Tensor Float 32
* Gaussian complex GEMMs using 3m complex multiply algorithm
* Universal GEMM kernel supporting two batch modes and two algorithms for parallel reductions
* Policy updates:
* [CUDA 11 Toolkit](https://developer.nvidia.com/cuda-toolkit) needed to enable NVIDIA Ampere Architecture features
* Disabled F16C by default for compatibility - enable on cmake command line with `-DCUTLASS_ENABLE_F16C=ON`
## [2.1.0](https://github.com/NVIDIA/cutlass/releases/tag/v2.1.0) (2020-04-06)
* BLAS-style host-side API added to [CUTLASS Library](/media/docs/quickstart.md#cutlass-library)
* API to launch compiled kernel instances for GEMM and planar complex GEMM
* Planar Complex GEMM kernels targeting Volta and Turing Tensor Cores
* Computes complex matrix products on matrices stored as disjoint real and imaginary parts
* [SDK Examples of Planar Complex GEMMs](/examples/10_planar_complex/planar_complex.cu)
* Minor enhancements and bug fixes
## [2.0.0](https://github.com/NVIDIA/cutlass/releases/tag/v2.0.0) (2019-11-19)
* Substantially refactored for
* Better performance, particularly for native Turing Tensor Cores
* Robust and durable templates spanning the design space
* Encapsulated functionality embodying modern C++11 programming techniques
* Optimized containers and data types for efficient, generic, portable device code
* Updates to:
* [Quick start guide](/media/docs/quickstart.md)
* [Documentation](/README.md#documentation)
* [Utilities](/media/docs/utilities.md)
* [CUTLASS Profiler](/media/docs/profiler.md)
* Native Turing Tensor Cores
* Efficient GEMM kernels targeting Turing Tensor Cores
* Mixed-precision floating point, 8-bit integer, 4-bit integer, and binarized operands
* Coverage of existing CUTLASS functionality
* GEMM kernels targeting CUDA and Tensor Cores in NVIDIA GPUs
* Volta Tensor Cores through native mma.sync and through WMMA API
* Optimizations such as parallel reductions, threadblock rasterization, and intra-threadblock reductions
* Batched GEMM operations
* Complex-valued GEMMs
* **Note: a host compiler supporting C++11 or greater is required.**
# CUTLASS 1.x
## [1.3.2](https://github.com/NVIDIA/cutlass/releases/tag/v1.3.2) (2019-07-09)
* Performance improvement for Volta Tensor Cores TN and TT layouts.
## [1.3.1](https://github.com/NVIDIA/cutlass/releases/tag/v1.3.1) (2019-04-09)
* Corrected NVRTC unit tests.
## [1.3.0](https://github.com/NVIDIA/cutlass/releases/tag/v1.3.0) (2019-03-20)
* Efficient GEMM kernel targeting Volta Tensor Cores via `mma.sync` instruction added in CUDA 10.1.
## [1.2.0](https://github.com/NVIDIA/cutlass/releases/tag/v1.2.0) (2018-10-26)
* Parallelized reductions across threadblocks ("Split-K")
* Improved IGEMM performance
* Batched strided WMMA GEMMs
## [1.1.0](https://github.com/NVIDIA/cutlass/releases/tag/v1.1.0) (2018-09-19)
* Turing Features
* WMMA GEMM targeting TensorCores - INT8, INT4, 1-bit
* Batched Strided GEMM
* Threadblock rasterization strategies
* Improved performance for adverse problem sizes and data layouts
* Extended CUTLASS Core comonents
* Tensor views support arbitrary matrix and tensor layouts
* Zip iterators for structuring multiple data streams
* Enhanced CUTLASS utilities
* Reference code for tensor operations in host and device code
* Added HostMatrix<> for simplified matrix creation
* Examples
* Basic GEMM, tensor views, CUTLASS utilities, batched GEMM, WMMA GEMM
## [1.0.1](https://github.com/NVIDIA/cutlass/releases/tag/v1.0.1) (2018-06-11)
* Intra-threadblock reduction added for small threadblock tile sizes
* sgemm_64x128x16, sgemm_128x128x16, sgemm_128x64x16, sgemm_128x32x16, sgemm_64x64x16, sgemm_64x32x16
* igemm_32x32x128
* GEMM _K_ residue handled during prologue prior to mainloop
* Replaced Google Test copy with submodule. Use `git submodule init --recursive --update`
## [1.0.0](https://github.com/NVIDIA/cutlass/commit/2028ebe120aab22bfd0b2baf8902d4c9627eb33f) (2018-05-16)
* Substantial rewrite to accommodate new architecture
* Kernels: SGEMM, DGEMM, IGEMM, HGEMM, WMMA GEMM
* Unit and performance tests
## [0.0.1](https://github.com/NVIDIA/cutlass/commit/d08ba8ac46e2fa3f745e070c390182edb56b2e91) (2017-12-04)
* Initial release
## Copyright
Copyright (c) 2017 - 2023 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
SPDX-License-Identifier: BSD-3-Clause
```
Redistribution and use in source and binary forms, with or without
modification, are permitted provided that the following conditions are met:
1. Redistributions of source code must retain the above copyright notice, this
list of conditions and the following disclaimer.
2. Redistributions in binary form must reproduce the above copyright notice,
this list of conditions and the following disclaimer in the documentation
and/or other materials provided with the distribution.
3. Neither the name of the copyright holder nor the names of its
contributors may be used to endorse or promote products derived from
this software without specific prior written permission.
THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
```

112
CITATION.cff Normal file
View File

@ -0,0 +1,112 @@
cff-version: 1.2.0
title: CUTLASS
message: >-
If you use this software, please cite using the
following metadata.
type: software
authors:
- given-names: Vijay
family-names: Thakkar
email: vithakkar@nvidia.com
affiliation: NVIDIA
- given-names: Pradeep
family-names: Ramani
email: prramani@nvidia.com
affiliation: NVIDIA
- given-names: Cris
family-names: Cecka
email: ccecka@nvidia.com
affiliation: NVIDIA
- given-names: Aniket
family-names: Shivam
email: ashivam@nvidia.com
affiliation: NVIDIA
- given-names: Honghao
family-names: Lu
email: honghaol@nvidia.com
affiliation: NVIDIA
- given-names: Ethan
family-names: Yan
email: etyan@nvidia.com
affiliation: NVIDIA
- given-names: Jack
family-names: Kosaian
email: jkosaian@nvidia.com
affiliation: NVIDIA
- given-names: Mark
family-names: Hoemmen
email: mhoemmen@nvidia.com
affiliation: NVIDIA
- given-names: Haicheng
family-names: Wu
email: haichengw@nvidia.com
affiliation: NVIDIA
- given-names: Andrew
family-names: Kerr
email: akerr@nvidia.com
affiliation: NVIDIA
- given-names: Matt
family-names: Nicely
email: mnicely@nvidia.com
affiliation: NVIDIA
- given-names: Duane
family-names: Merrill
email: dumerrill@nvidia.com
affiliation: NVIDIA
- given-names: Dustyn
family-names: Blasig
email: dblasig@nvidia.com
affiliation: NVIDIA
- given-names: Fengqi
family-names: Qiao
email: fqiao@nvidia.com
affiliation: NVIDIA
- given-names: Piotr
family-names: Majcher
email: pmajcher@nvidia.com
affiliation: NVIDIA
- given-names: Paul
family-names: Springer
email: pspringer@nvidia.com
affiliation: NVIDIA
- given-names: Markus
family-names: Hohnerbach
affiliation: NVIDIA
email: mhohnerbach@nvidia.com
- given-names: Jin
family-names: Wang
email: jinw@nvidia.com
affiliation: NVIDIA
- given-names: Manish
family-names: Gupta
affiliation: Google
email: manigupta@google.com
repository-code: 'https://github.com/NVIDIA/cutlass'
abstract: >-
CUTLASS is a collection of CUDA C++ template
abstractions for implementing high-performance
matrix-multiplication (GEMM) and related
computations at all levels and scales within CUDA.
It incorporates strategies for hierarchical
decomposition and data movement similar to those
used to implement cuBLAS and cuDNN. CUTLASS
decomposes these "moving parts" into reusable,
modular software components abstracted by C++
template classes. These thread-wide, warp-wide,
block-wide, and device-wide primitives can be
specialized and tuned via custom tiling sizes, data
types, and other algorithmic policy. The resulting
flexibility simplifies their use as building blocks
within custom kernels and applications.
keywords:
- 'cutlass, tensor cores, cuda, cute, nvidia, gpu, linear algebra, matrix computations'
license: BSD-3-Clause
license-url: https://github.com/NVIDIA/cutlass/blob/v3.0.0/LICENSE.txt
version: '3.0.0'
date-released: '2023-01-23'
identifiers:
- type: url
value: "https://github.com/NVIDIA/cutlass/tree/v3.0.0"
description: The GitHub release URL of tag 3.0.0

840
CMakeLists.txt Executable file
View File

@ -0,0 +1,840 @@
# Copyright (c) 2017 - 2023 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: BSD-3-Clause
#
# Redistribution and use in source and binary forms, with or without
# modification, are permitted provided that the following conditions are met:
#
# 1. Redistributions of source code must retain the above copyright notice, this
# list of conditions and the following disclaimer.
#
# 2. Redistributions in binary form must reproduce the above copyright notice,
# this list of conditions and the following disclaimer in the documentation
# and/or other materials provided with the distribution.
#
# 3. Neither the name of the copyright holder nor the names of its
# contributors may be used to endorse or promote products derived from
# this software without specific prior written permission.
#
# THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
# AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
# IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
# DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
# FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
# DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
# SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
# CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
# OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
# OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
cmake_minimum_required(VERSION 3.18 FATAL_ERROR)
if(cutlass_LOADED)
# If CUTLASS has been previously fetched and loaded, don't do it again.
return()
else()
set(cutlass_LOADED ON)
set(CUTLASS_DIR ${CMAKE_CURRENT_SOURCE_DIR} CACHE PATH "CUTLASS Repository Directory")
endif()
message(STATUS "CMake Version: ${CMAKE_VERSION}")
set(IMPLICIT_CMAKE_CXX_STANDARD OFF CACHE BOOL "Do not explicitly specify -std=c++11 if set")
project(CUTLASS VERSION 3.0.0 LANGUAGES CXX)
include(${CMAKE_CURRENT_SOURCE_DIR}/CUDA.cmake)
if (CUDA_VERSION VERSION_LESS 11.3)
message(WARNING "CUTLASS ${CUTLASS_VERSION} requires CUDA 11.4 or higher, and strongly recommends CUDA 11.8 or higher.")
elseif (CUDA_VERSION VERSION_LESS 11.4)
message(WARNING "CUTLASS ${CUTLASS_VERSION} support for CUDA ${CUDA_VERSION} is deprecated, please use CUDA 11.8 or higher.")
endif()
if(CMAKE_CXX_COMPILER_ID STREQUAL "GNU" AND CMAKE_CXX_COMPILER_VERSION VERSION_LESS 7.5)
message(FATAL_ERROR "GCC version must be at least 7.5!")
endif()
if (CUDA_COMPILER MATCHES "[Cc]lang" AND CMAKE_CXX_COMPILER_VERSION VERSION_LESS 7.0)
message(FATAL_ERROR "Clang 7.0+ required for GPU compilation")
endif()
find_package(Doxygen QUIET)
#
# CUTLASS 3.x requires C++17
#
set(CMAKE_CXX_STANDARD 17)
set(CMAKE_CXX_STANDARD_REQUIRED ON)
set(CMAKE_CXX_EXTENSIONS OFF)
if(CUTLASS_NATIVE_CUDA)
set(CMAKE_CUDA_STANDARD 17)
set(CMAKE_CUDA_STANDARD_REQUIRED ON)
list(APPEND CUTLASS_CUDA_NVCC_FLAGS --expt-relaxed-constexpr)
else()
list(APPEND CUTLASS_CUDA_NVCC_FLAGS --std=c++17)
endif()
if(CMAKE_INSTALL_PREFIX_INITIALIZED_TO_DEFAULT)
set(CMAKE_INSTALL_PREFIX install CACHE PATH "Default installation location." FORCE)
endif()
message(STATUS "Default Install Location: ${CMAKE_INSTALL_PREFIX}")
set(CUTLASS_ENABLE_HEADERS_ONLY OFF CACHE BOOL "Enable only the header library")
if(CUTLASS_ENABLE_HEADERS_ONLY)
set(CUTLASS_ENABLE_EXAMPLES_INIT OFF)
set(CUTLASS_ENABLE_TOOLS_INIT ON)
set(CUTLASS_ENABLE_LIBRARY_INIT OFF)
else()
set(CUTLASS_ENABLE_EXAMPLES_INIT ON)
set(CUTLASS_ENABLE_TOOLS_INIT ON)
set(CUTLASS_ENABLE_LIBRARY_INIT ON)
endif()
set(CUTLASS_TEST_UNIT_ENABLE_WARNINGS OFF CACHE BOOL "Enable warnings on waived unit tests.")
set(CUTLASS_ENABLE_EXAMPLES ${CUTLASS_ENABLE_EXAMPLES_INIT} CACHE BOOL "Enable CUTLASS Examples")
set(CUTLASS_ENABLE_TOOLS ${CUTLASS_ENABLE_TOOLS_INIT} CACHE BOOL "Enable CUTLASS Tools")
set(CUTLASS_ENABLE_LIBRARY ${CUTLASS_ENABLE_LIBRARY_INIT} CACHE BOOL "Enable CUTLASS Library")
set(CUTLASS_ENABLE_PROFILER ${CUTLASS_ENABLE_LIBRARY} CACHE BOOL "Enable CUTLASS Profiler")
set(CUTLASS_ENABLE_PERFORMANCE ${CUTLASS_ENABLE_PROFILER} CACHE BOOL "Enable CUTLASS Proformance")
if(${CMAKE_PROJECT_NAME} STREQUAL ${PROJECT_NAME})
set(CUTLASS_ENABLE_TESTS_INIT ${CUTLASS_ENABLE_LIBRARY}})
else()
set(CUTLASS_ENABLE_TESTS_INIT OFF)
endif()
set(CUTLASS_ENABLE_TESTS ${CUTLASS_ENABLE_TESTS_INIT} CACHE BOOL "Enable CUTLASS Tests")
if (CUTLASS_ENABLE_TESTS)
include(${CMAKE_CURRENT_SOURCE_DIR}/cmake/googletest.cmake)
endif()
set(CUTLASS_NVCC_ARCHS_SUPPORTED "")
if (CUDA_VERSION VERSION_GREATER_EQUAL 11.4 AND NOT CUDA_COMPILER MATCHES "[Cc]lang")
list(APPEND CUTLASS_NVCC_ARCHS_SUPPORTED 70 72 75 80 86 87)
endif()
if (CUDA_VERSION VERSION_GREATER_EQUAL 11.8 AND NOT CUDA_COMPILER MATCHES "[Cc]lang")
list(APPEND CUTLASS_NVCC_ARCHS_SUPPORTED 89 90)
endif()
if (CUDA_VERSION VERSION_GREATER_EQUAL 12.0 AND NOT CUDA_COMPILER MATCHES "[Cc]lang")
list(APPEND CUTLASS_NVCC_ARCHS_SUPPORTED 90a)
endif()
set(CUTLASS_NVCC_ARCHS ${CUTLASS_NVCC_ARCHS_SUPPORTED} CACHE STRING "The SM architectures requested.")
set(CUTLASS_NVCC_ARCHS_ENABLED ${CUTLASS_NVCC_ARCHS} CACHE STRING "The SM architectures to build code for.")
# Special policy introduced in CMake 3.13
if (POLICY CMP0076)
cmake_policy(SET CMP0076 NEW)
endif()
include(GNUInstallDirs)
link_directories(${CUDA_TOOLKIT_ROOT_DIR}/lib64/stubs)
###################################################################################################
#
# Configure CMake variables
#
###################################################################################################
message(STATUS "CUDA Compilation Architectures: ${CUTLASS_NVCC_ARCHS_ENABLED}")
if (NOT (CMAKE_BUILD_TYPE OR CONFIGURATION_TYPES))
# By default we want to build in Release mode to ensure that we're getting best performance.
set(CMAKE_BUILD_TYPE Release CACHE STRING "Choose build level" FORCE)
set_property(CACHE CMAKE_BUILD_TYPE PROPERTY STRINGS "Debug" "RelWithDebInfo" "Release")
endif()
set(CMAKE_POSITION_INDEPENDENT_CODE ON)
if (DEFINED CMAKE_DEBUG_POSTFIX)
set(CUTLASS_LIBRARY_DEBUG_POSTFIX_INIT ${CMAKE_DEBUG_POSTFIX})
else()
set(CUTLASS_LIBRARY_DEBUG_POSTFIX_INIT .debug)
endif()
set(CUTLASS_LIBRARY_DEBUG_POSTFIX ${CUTLASS_LIBRARY_DEBUG_POSTFIX_INIT} CACHE STRING "Default postfix value for debug libraries")
if(WIN32)
# On Windows we link against the shared (DLL) runtime. Change gtest settings to match this.
set(gtest_force_shared_crt ON CACHE BOOL "Use shared (DLL) run-time lib even when Google Test is built as static lib" FORCE)
endif()
if (WIN32)
# Enable more warnings and treat as errors
list(APPEND CUTLASS_CUDA_NVCC_FLAGS -Xcompiler=/W3 -Xcompiler=/WX)
# Disable warning on Unicode characters
list(APPEND CUTLASS_CUDA_NVCC_FLAGS -Xcompiler=/wd4819)
# Disable excess x86 floating point precision that can lead to results being labeled incorrectly
list(APPEND CUTLASS_CUDA_NVCC_FLAGS -Xcompiler=/fp:strict)
endif(WIN32)
if (${CUTLASS_NVCC_VERBOSE})
list(APPEND CUTLASS_CUDA_NVCC_FLAGS -v)
endif()
#
# CUTLASS NAMESPACE
#
set(CUTLASS_NAMESPACE "cutlass" CACHE STRING "Top level namespace of CUTLASS")
set(CUTLASS_NVCC_EMBED_CUBIN ON CACHE BOOL "Embed compiled CUDA kernel binaries into executables.")
set(CUTLASS_NVCC_EMBED_PTX ON CACHE BOOL "Embed compiled PTX into executables.")
set(CUTLASS_NVCC_KEEP OFF CACHE BOOL "Keep intermediate files generated by NVCC.")
set(CUTLASS_ENABLE_F16C OFF CACHE BOOL "Enable F16C x86 extensions in host code.")
#
# CUTLASS generator cmake configuration
#
set(CUTLASS_LIBRARY_OPERATIONS "all" CACHE STRING "Comma delimited list of operation name filters. Default '' means all operations are enabled.")
set(CUTLASS_LIBRARY_KERNELS "" CACHE STRING "Comma delimited list of kernel name filters. If unspecified, only the largest tile size is enabled. If 'all' is specified, all kernels are enabled.")
set(CUTLASS_LIBRARY_IGNORE_KERNELS "" CACHE STRING "Comma delimited list of kernel names to exclude from build.")
# Test Levels L0, L1, L2
set(CUTLASS_TEST_LEVEL "0" CACHE STRING "Level of tests to compile.")
set(CUTLASS_TEST_ENABLE_CACHED_RESULTS ON CACHE BOOL "Enable caching and reuse of test results in unit tests")
set_property(CACHE CUTLASS_TEST_LEVEL PROPERTY STRINGS 0 1 2)
list(APPEND CUTLASS_CUDA_NVCC_FLAGS -DCUTLASS_TEST_LEVEL=${CUTLASS_TEST_LEVEL})
list(APPEND CUTLASS_CUDA_CLANG_FLAGS -DCUTLASS_TEST_LEVEL=${CUTLASS_TEST_LEVEL})
if (CUTLASS_TEST_ENABLE_CACHED_RESULTS)
message(STATUS "Enable caching of reference results in conv unit tests")
list(APPEND CUTLASS_CUDA_NVCC_FLAGS -DCUTLASS_TEST_ENABLE_CACHED_RESULTS=1)
endif()
set(CUTLASS_CONV_UNIT_TEST_RIGOROUS_SIZE_ENABLED ON CACHE BOOL "Enable/Disable rigorous conv problem sizes in conv unit tests")
if (CUTLASS_CONV_UNIT_TEST_RIGOROUS_SIZE_ENABLED)
message(STATUS "Enable rigorous conv problem sizes in conv unit tests")
list(APPEND CUTLASS_CUDA_NVCC_FLAGS -DCUTLASS_CONV_UNIT_TEST_RIGOROUS_SIZE_ENABLED=1)
endif()
#
# CUDA 10.1 introduces "mma" in PTX performing collective matrix multiply operations.
#
if (CUDA_VERSION VERSION_LESS 10.1)
set(CUTLASS_ENABLE_TENSOR_CORE_MMA_DEFAULT OFF)
else()
set(CUTLASS_ENABLE_TENSOR_CORE_MMA_DEFAULT ON)
endif()
# Trace levels for debugging
set(CUTLASS_DEBUG_TRACE_LEVEL "0" CACHE STRING "Level of debug tracing to perform.")
list(APPEND CUTLASS_CUDA_NVCC_FLAGS -DCUTLASS_DEBUG_TRACE_LEVEL=${CUTLASS_DEBUG_TRACE_LEVEL})
set(CUTLASS_ENABLE_TENSOR_CORE_MMA ${CUTLASS_ENABLE_TENSOR_CORE_MMA_DEFAULT} CACHE BOOL
"Enable PTX mma instruction for collective matrix multiply operations.")
#
# NOTE: running with asan and CUDA requires the following environment variable:
#
# ASAN_OPTIONS=protect_shadow_gap=0:replace_intrin=0:detect_leaks=0
#
# without the above environment setting, an error like the following may be generated:
#
# *** Error: Could not detect active GPU device ID [out of memory]
# ...
# ==9149==ERROR: LeakSanitizer: detected memory leaks
# ...
#
if(ENABLE_ASAN) # https://github.com/google/sanitizers/wiki/AddressSanitizer
list(APPEND CUTLASS_CUDA_NVCC_FLAGS --compiler-options=-fsanitize=address --compiler-options=-fno-omit-frame-pointer)
string(APPEND CMAKE_EXE_LINKER_FLAGS " -fsanitize=address")
endif()
###################################################################################################
#
# Configure CUDA build options
#
###################################################################################################
if(CUTLASS_NVCC_EMBED_PTX)
list(APPEND CUTLASS_CUDA_CLANG_FLAGS --cuda-include-ptx=all)
endif()
if (CUTLASS_ENABLE_TENSOR_CORE_MMA)
list(APPEND CUTLASS_CUDA_FLAGS -DCUTLASS_ENABLE_TENSOR_CORE_MMA=1)
endif()
if (NOT MSVC AND CUTLASS_NVCC_KEEP)
# MSVC flow handles caching already, but for other generators we handle it here.
set(CUTLASS_NVCC_KEEP_DIR ${CMAKE_CURRENT_BINARY_DIR}/tmp CACHE PATH "Location to store NVCC scratch files")
file(MAKE_DIRECTORY ${CUTLASS_NVCC_KEEP_DIR})
list(APPEND CUTLASS_CUDA_NVCC_FLAGS --keep -v) # --keep-dir may not work with nvcc for some directories.
list(APPEND CUTLASS_CUDA_CLANG_FLAGS -save-temps=${CUTLASS_NVCC_KEEP_DIR})
endif()
if (CUTLASS_ENABLE_F16C AND NOT CMAKE_CROSSCOMPILING)
list(APPEND CUTLASS_CUDA_FLAGS -DCUTLASS_ENABLE_F16C=1)
if ((CMAKE_CXX_COMPILER_ID MATCHES "GNU") OR (CMAKE_CXX_COMPILER_ID MATCHES "Clang"))
list(APPEND CUTLASS_CUDA_NVCC_FLAGS -Xcompiler=-mf16c)
elseif((CMAKE_CXX_COMPILER_ID MATCHES "MSVC"))
list(APPEND CUTLASS_CUDA_NVCC_FLAGS -Xcompiler=/arch:AVX2)
endif()
endif()
if (CUTLASS_ENABLE_OPENMP_TESTS)
find_package(OpenMP)
if(OpenMP_CXX_FOUND)
list(APPEND CUTLASS_CUDA_NVCC_FLAGS -Xcompiler=${OpenMP_CXX_FLAGS})
else()
message(WARNING "CUTLASS_ENABLE_OPENMP_TESTS set but OpenMP not found.")
endif()
endif()
list(APPEND CUTLASS_CUDA_NVCC_FLAGS $<$<BOOL:${UNIX}>:-Xcompiler=-Wconversion>)
list(APPEND CUTLASS_CUDA_NVCC_FLAGS $<$<BOOL:${UNIX}>:-Xcompiler=-fno-strict-aliasing>)
# Don't leak lineinfo in release builds
if (NOT CMAKE_BUILD_TYPE MATCHES "Release")
list(APPEND CUTLASS_CUDA_CLANG_FLAGS -gmlt)
list(APPEND CUTLASS_CUDA_NVCC_FLAGS -lineinfo)
endif()
#Report CUDA build flags
if (CUDA_COMPILER MATCHES "[Cc]lang")
if(CUTLASS_CUDA_CLANG_FLAGS)
message(STATUS "Using CLANG flags: ${CUTLASS_CUDA_CLANG_FLAGS}")
endif()
else()
if(CUTLASS_CUDA_NVCC_FLAGS)
message(STATUS "Using NVCC flags: ${CUTLASS_CUDA_NVCC_FLAGS}")
endif()
endif()
if(CUDA_COMPILER MATCHES "[Cc]lang")
if( NOT CMAKE_CXX_COMPILER_ID MATCHES "Clang" )
message(FATAL_ERROR "Clang CUDA compilation requires Clang CXX compilation. Currently CMAKE_CXX_COMPILER is ${CMAKE_CXX_COMPILER_ID}" )
endif()
# There are numerous Clang versions that can work with each CUDA toolkit and the
# the checks are not very useful so we are turning them off and using testing to
# ensure the various combinations work properly.
list(APPEND CUTLASS_CUDA_CLANG_FLAGS --cuda-path=${CUDA_TOOLKIT_ROOT_DIR})
list(APPEND CUTLASS_CUDA_CLANG_FLAGS -D__NV_NO_HOST_COMPILER_CHECK=1)
list(APPEND CUTLASS_CUDA_CLANG_FLAGS -Wno-unknown-cuda-version)
list(APPEND CUTLASS_CUDA_CLANG_FLAGS -mllvm -pragma-unroll-threshold=100000)
list(APPEND CUTLASS_CUDA_CLANG_FLAGS -mllvm -unroll-threshold=5000)
list(APPEND CUTLASS_CUDA_CLANG_FLAGS -Wno-unused-command-line-argument)
string(REPLACE "." ";" CUDA_VERSION_PARTS ${CMAKE_CUDA_COMPILER_VERSION})
list(GET CUDA_VERSION_PARTS 0 CUDA_VERSION_MAJOR)
list(GET CUDA_VERSION_PARTS 1 CUDA_VERSION_MINOR)
list(APPEND CUTLASS_CUDA_CLANG_FLAGS -D__CUDACC_VER_MAJOR__=${CUDA_VERSION_MAJOR} -D__CUDACC_VER_MINOR__=${CUDA_VERSION_MINOR})
# needed for libcublasLt.so in case it's installed in the same location as libcudart.so
# dynamic linker can find it if linker sets RPATH (forced by --disable-new-tags)
# Otherwise linker uses RUNPATH and that does not propagate to loaded libs.
list(APPEND CUTLASS_CUDA_CLANG_FLAGS -Wl,--disable-new-dtags)
link_libraries(nvidia::cudart)
link_libraries(nvidia::cuda_driver)
endif()
# Support for 128-bit integers if using NVIDIA C++ compiler
if (${CMAKE_CXX_COMPILER_ID} MATCHES "PGI" OR ${CMAKE_CXX_COMPILER_ID} MATCHES "NVHPC")
set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -Mint128 ")
endif()
if (CMAKE_VERSION VERSION_GREATER_EQUAL 3.18)
# CMake 3.18 added support for CUDA_ARCHITECTURES target property. We will use this
# property for CMake 3.18+, so we request the NEW behavior for correct compatibility.
# https://cmake.org/cmake/help/v3.18/policy/CMP0104.html#policy:CMP0104
cmake_policy(SET CMP0104 NEW)
endif()
function(cutlass_apply_cuda_gencode_flags TARGET)
set(options)
set(oneValueArgs)
set(multiValueArgs SM_ARCHS)
cmake_parse_arguments(_ "${options}" "${oneValueArgs}" "${multiValueArgs}" ${ARGN})
if (__SM_ARCHS)
set(ARCHS_ENABLED ${__SM_ARCHS})
else()
set(ARCHS_ENABLED ${CUTLASS_NVCC_ARCHS_ENABLED})
endif()
set(NVCC_FLAGS)
set(CLANG_FLAGS)
set(__CMAKE_CUDA_ARCHS)
foreach(ARCH ${ARCHS_ENABLED})
list(APPEND CLANG_FLAGS --cuda-gpu-arch=sm_${ARCH})
set(CODES)
if(CUTLASS_NVCC_EMBED_CUBIN)
list(APPEND CODES sm_${ARCH})
list(APPEND __CMAKE_CUDA_ARCHS ${ARCH}-real)
endif()
if(CUTLASS_NVCC_EMBED_PTX)
list(APPEND CODES compute_${ARCH})
list(APPEND __CMAKE_CUDA_ARCHS ${ARCH}-virtual)
endif()
list(JOIN CODES "," CODES_STR)
list(APPEND NVCC_FLAGS -gencode=arch=compute_${ARCH},code=[${CODES_STR}])
endforeach()
if (NOT __SM_ARCHS)
if (CUDA_COMPILER MATCHES "[Cc]lang")
target_compile_options(
${TARGET}
PRIVATE
$<$<COMPILE_LANGUAGE:CXX>:${CLANG_FLAGS}>
)
elseif(CMAKE_VERSION GREATER_EQUAL 3.18)
set_property(TARGET ${TARGET} PROPERTY CUDA_ARCHITECTURES ${__CMAKE_CUDA_ARCHS})
else()
target_compile_options(
${TARGET}
PRIVATE
$<$<COMPILE_LANGUAGE:CUDA>:${NVCC_FLAGS}>
)
endif()
else()
list(JOIN CLANG_FLAGS " " CLANG_FLAGS_STR)
list(JOIN NVCC_FLAGS " " STR_NVCC_FLAGS)
if (CUDA_COMPILER MATCHES "[Cc]lang")
if(${TARGET} MATCHES ".*\.cpp")
set_source_files_properties(${TARGET} PROPERTIES COMPILE_FLAGS ${CLANG_FLAGS_STR})
endif()
elseif(CMAKE_VERSION GREATER_EQUAL 3.18)
set_source_files_properties(${TARGET} PROPERTIES CUDA_ARCHITECTURES ${STR_NVCC_FLAGS})
else()
if(${TARGET} MATCHES ".*\.cu")
set_source_files_properties(${TARGET} PROPERTIES COMPILE_FLAGS ${STR_NVCC_FLAGS})
endif()
endif()
endif()
endfunction()
# Cache the flags so they are available when the function below is called anywhere globally.
set(__CUTLASS_CUDA_FLAGS ${CUTLASS_CUDA_FLAGS} CACHE INTERNAL "")
set(__CUTLASS_CUDA_FLAGS_RELEASE ${CUTLASS_CUDA_FLAGS_RELEASE} CACHE INTERNAL "")
set(__CUTLASS_CUDA_FLAGS_RELWITHDEBINFO ${CUTLASS_CUDA_FLAGS_RELWITHDEBINFO} CACHE INTERNAL "")
set(__CUTLASS_CUDA_FLAGS_DEBUG ${CUTLASS_CUDA_FLAGS_DEBUG} CACHE INTERNAL "")
set(__CUTLASS_CUDA_CLANG_FLAGS ${CUTLASS_CUDA_CLANG_FLAGS} CACHE INTERNAL "")
set(__CUTLASS_CUDA_CLANG_FLAGS_RELEASE ${CUTLASS_CUDA_CLANG_FLAGS_RELEASE} CACHE INTERNAL "")
set(__CUTLASS_CUDA_CLANG_FLAGS_RELWITHDEBINFO ${CUTLASS_CUDA_CLANG_FLAGS_RELWITHDEBINFO} CACHE INTERNAL "")
set(__CUTLASS_CUDA_CLANG_FLAGS_DEBUG ${CUTLASS_CUDA_CLANG_FLAGS_DEBUG} CACHE INTERNAL "")
set(__CUTLASS_CUDA_NVCC_FLAGS ${CUTLASS_CUDA_NVCC_FLAGS} CACHE INTERNAL "")
set(__CUTLASS_CUDA_NVCC_FLAGS_RELEASE ${CUTLASS_CUDA_NVCC_FLAGS_RELEASE} CACHE INTERNAL "")
set(__CUTLASS_CUDA_NVCC_FLAGS_RELWITHDEBINFO ${CUTLASS_CUDA_NVCC_FLAGS_RELWITHDEBINFO} CACHE INTERNAL "")
set(__CUTLASS_CUDA_NVCC_FLAGS_DEBUG ${CUTLASS_CUDA_NVCC_FLAGS_DEBUG} CACHE INTERNAL "")
function(cutlass_apply_standard_compile_options TARGET)
if(CUDA_COMPILER MATCHES "[Cc]lang")
set(CUDA_COMPILE_LANGUAGE CXX)
set(_FLAGS ${__CUTLASS_CUDA_FLAGS} ${__CUTLASS_CUDA_CLANG_FLAGS})
set(_FLAGS_RELEASE ${__CUTLASS_CUDA_FLAGS_RELEASE} ${__CUTLASS_CUDA_CLANG_FLAGS_RELEASE})
set(_FLAGS_RELWITHDEBINFO ${__CUTLASS_CUDA_FLAGS_RELWITHDEBINFO} ${__CUTLASS_CUDA_CLANG_FLAGS_RELWITHDEBINFO})
set(_FLAGS_DEBUG ${__CUTLASS_CUDA_FLAGS_DEBUG} ${__CUTLASS_CUDA_CLANG_FLAGS_DEBUG})
else()
set(CUDA_COMPILE_LANGUAGE CUDA)
set(_FLAGS ${__CUTLASS_CUDA_FLAGS} ${__CUTLASS_CUDA_NVCC_FLAGS})
set(_FLAGS_RELEASE ${__CUTLASS_CUDA_FLAGS_RELEASE} ${__CUTLASS_CUDA_NVCC_FLAGS_RELEASE})
set(_FLAGS_RELWITHDEBINFO ${__CUTLASS_CUDA_FLAGS_RELWITHDEBINFO} ${__CUTLASS_CUDA_NVCC_FLAGS_RELWITHDEBINFO})
set(_FLAGS_DEBUG ${__CUTLASS_CUDA_FLAGS_DEBUG} ${__CUTLASS_CUDA_NVCC_FLAGS_DEBUG})
endif()
target_link_libraries(${TARGET} PRIVATE CUTLASS)
target_compile_options(
${TARGET}
PRIVATE
$<$<COMPILE_LANGUAGE:${CUDA_COMPILE_LANGUAGE}>:${_FLAGS}>
$<$<COMPILE_LANGUAGE:${CUDA_COMPILE_LANGUAGE}>:$<$<CONFIG:RELEASE>:${_FLAGS_RELEASE}>>
$<$<COMPILE_LANGUAGE:${CUDA_COMPILE_LANGUAGE}>:$<$<CONFIG:RELWITHDEBINFO>:${_FLAGS_RELWITHDEBINFO}>>
$<$<COMPILE_LANGUAGE:${CUDA_COMPILE_LANGUAGE}>:$<$<CONFIG:DEBUG>:${_FLAGS_DEBUG}>>
)
endfunction()
#
# The following items should eventually be pushed into cutlass/CMakeLists.txt
#
# GLOB for CUTLASS header files. Should we use a static list instead?
file(GLOB_RECURSE CUTLASS_INCLUDE RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} include/cutlass/*.h)
file(GLOB_RECURSE CUTLASS_CUTLASS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR}/include include/cutlass/*.h)
file(GLOB_RECURSE CUTLASS_NVRTC RELATIVE ${CMAKE_CURRENT_SOURCE_DIR}/test test/unit/nvrtc/kernel/*.h)
###################################################################################################
#
# Define build targets
#
###################################################################################################
source_group(TREE ${CMAKE_CURRENT_SOURCE_DIR}/include REGULAR_EXPRESSION ".*\.h")
add_library(CUTLASS INTERFACE)
add_library(nvidia::cutlass::cutlass ALIAS CUTLASS)
set_target_properties(CUTLASS PROPERTIES EXPORT_NAME cutlass)
set(CUTLASS_INCLUDE_DIR ${CMAKE_CURRENT_SOURCE_DIR}/include CACHE PATH "CUTLASS Header Library")
set(CUTLASS_GENERATOR_DIR ${CMAKE_CURRENT_SOURCE_DIR}/tools/library CACHE INTERNAL "Location of generator scripts")
# The following utility directory is needed even if the tools build is disabled, so it exists here.
set(CUTLASS_TOOLS_UTIL_INCLUDE_DIR ${CMAKE_CURRENT_SOURCE_DIR}/tools/util/include CACHE INTERNAL "")
include_directories(${CUTLASS_INCLUDE_DIR})
target_compile_features(CUTLASS INTERFACE cxx_std_11)
if (NOT CUTLASS_NAMESPACE STREQUAL "cutlass")
target_compile_definitions(CUTLASS INTERFACE CUTLASS_NAMESPACE=${CUTLASS_NAMESPACE})
endif()
if (NOT DEFINED CUTLASS_REVISION)
find_package(Git QUIET)
execute_process(
COMMAND ${GIT_EXECUTABLE} rev-parse --short HEAD
RESULT_VARIABLE CUTLASS_REVISION_RESULT
OUTPUT_VARIABLE CUTLASS_REVISION
OUTPUT_STRIP_TRAILING_WHITESPACE
)
if (CUTLASS_REVISION_RESULT)
message(STATUS "CUTLASS Revision: Unable to detect, Git returned code ${CUTLASS_REVISION_RESULT}.")
else()
message(STATUS "CUTLASS Revision: ${CUTLASS_REVISION}")
endif()
endif()
configure_file(
${CMAKE_CURRENT_SOURCE_DIR}/cmake/version.h.in
${CMAKE_CURRENT_BINARY_DIR}/include/cutlass/version.h
@ONLY)
target_include_directories(
CUTLASS
INTERFACE
$<INSTALL_INTERFACE:include>
$<BUILD_INTERFACE:${CUTLASS_INCLUDE_DIR}>
$<BUILD_INTERFACE:${CMAKE_CURRENT_BINARY_DIR}/include>
$<BUILD_INTERFACE:${CUDA_TOOLKIT_ROOT_DIR}/include>
$<BUILD_INTERFACE:${cute_SOURCE_DIR}/include>
$<BUILD_INTERFACE:${cute_SOURCE_DIR}/examples>
)
install(
DIRECTORY
${CUTLASS_INCLUDE_DIR}/
${CMAKE_CURRENT_BINARY_DIR}/include/
DESTINATION ${CMAKE_INSTALL_INCLUDEDIR}
)
install(
TARGETS CUTLASS
EXPORT NvidiaCutlass
PUBLIC_HEADER DESTINATION ${CMAKE_INSTALL_INCLUDEDIR}
)
################################################################################
# Doxygen is available. Generate documentation
if (DOXYGEN_FOUND)
# DOT is available. Enable graph generation in the documentation
if (DOXYGEN_DOT_EXECUTABLE)
set(CUTLASS_ENABLE_DOXYGEN_DOT ON CACHE BOOL "Use dot to generate graphs in the doxygen documentation.")
else()
set(CUTLASS_ENABLE_DOXYGEN_DOT OFF CACHE BOOL "Use dot to generate graphs in the doxygen documentation." FORCE)
endif()
if (CUTLASS_ENABLE_DOXYGEN_DOT)
set(HAVE_DOT "YES")
else()
set(HAVE_DOT "NO")
endif()
# Add custom target for Doxygen.
add_custom_target(cutlass_docs ${CMAKE_COMMAND} -E env
"DOT_PATH=${DOXYGEN_DOT_EXECUTABLE}"
"HAVE_DOT=${HAVE_DOT}"
${DOXYGEN_EXECUTABLE} ${CMAKE_CURRENT_SOURCE_DIR}/Doxyfile
WORKING_DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR}
VERBATIM
)
endif()
if(NOT WIN32)
# Add common library search paths so executables and libraries can load and run
# without LD_LIBRARY_PATH being set.
link_libraries(
"-Wl,-rpath,'$ORIGIN'"
"-Wl,-rpath,'$ORIGIN/../lib64'"
"-Wl,-rpath,'$ORIGIN/../lib'"
"-Wl,-rpath,'${CUDA_TOOLKIT_ROOT_DIR}/lib64'"
"-Wl,-rpath,'${CUDA_TOOLKIT_ROOT_DIR}/lib'"
)
endif()
################################################################################
include(CTest)
enable_testing()
if (NOT TARGET test_all)
add_custom_target(test_all)
endif()
set(CUTLASS_INSTALL_TESTS ON CACHE BOOL "Install test executables")
set(CUTLASS_TEST_EXECUTION_ENVIRONMENT "" CACHE BOOL "Environment in which to invoke unit test executables")
set(CMAKE_TEST_INSTALL_PREFIX test CACHE STRING "Test root install location, relative to CMAKE_INSTALL_PREFIX.")
set(CUTLASS_TEST_INSTALL_PREFIX ${CMAKE_TEST_INSTALL_PREFIX}/cutlass CACHE STRING "Test root install location, relative to CMAKE_INSTALL_PREFIX.")
set(CUTLASS_TEST_INSTALL_BINDIR ${CUTLASS_TEST_INSTALL_PREFIX}/${CMAKE_INSTALL_BINDIR} CACHE STRING "Test root install location, relative to CMAKE_INSTALL_PREFIX.")
set(CUTLASS_TEST_INSTALL_LIBDIR ${CUTLASS_TEST_INSTALL_PREFIX}/${CMAKE_INSTALL_LIBDIR} CACHE STRING "Test root install location, relative to CMAKE_INSTALL_PREFIX.")
install(DIRECTORY DESTINATION ${CUTLASS_TEST_INSTALL_PREFIX})
install(DIRECTORY DESTINATION ${CUTLASS_TEST_INSTALL_BINDIR})
install(DIRECTORY DESTINATION ${CUTLASS_TEST_INSTALL_LIBDIR})
install(DIRECTORY DESTINATION ${CUTLASS_TEST_INSTALL_PREFIX}/ctest)
################################################################################
set(CUTLASS_ENABLE_CUBLAS OFF CACHE BOOL "cuBLAS usage for tests")
set(CUTLASS_ENABLE_CUDNN OFF CACHE BOOL "cuDNN usage for tests")
include(${CMAKE_CURRENT_SOURCE_DIR}/cuBLAS.cmake)
if (CUTLASS_ENABLE_CUBLAS)
target_compile_definitions(CUTLASS INTERFACE CUTLASS_ENABLE_CUBLAS=1)
endif()
include(${CMAKE_CURRENT_SOURCE_DIR}/cuDNN.cmake)
if (CUTLASS_ENABLE_CUDNN)
target_compile_definitions(CUTLASS INTERFACE CUTLASS_ENABLE_CUDNN=1)
endif()
################################################################################
set(CUTLASS_CTEST_TEMPLATE_FILE ${CMAKE_CURRENT_LIST_DIR}/cmake/CTestTestfile.config.cmake)
set(CUTLASS_CTEST_GENERATED_FILES "" CACHE INTERNAL "")
function(cutlass_add_executable_tests NAME TARGET)
#
# Generates test rules for `make test`, `make test_all`, and `ctest` invoked from either the
# <CMAKE_BINARY_DIR> or the <CMAKE_INSTALL_PREFIX>/<CUTLASS_TEST_INSTALL_PREFIX> after installation.
#
# NAME: The base name for the test. Can be run with `make <NAME>` or `ctest -R 'c<NAME>'`.
# TARGET: The target corresponding to the executable under test.
# DISABLE_EXECUTABLE_INSTALL_RULE: An option, if given, that disables creating an install rule for TARGET.
# DEPENDS: A list of targets or files on which this test is dependent.
# DEPENDEES: A list of targets which should depend on this test.
# TEST_COMMAND_OPTIONS: A list of variables (i.e. by reference params) which contain command line arguments
# to pass to the test executable. A unique test with suffix _0, _1, ... is generated for each set of
# options given. If this option is not used, a single test with no arguments is generated.
# RESULT_CACHE_FILE: A file to be installed alongside the test executable with pre-computed
# test results to speed up test runtime.
#
set(options DISABLE_EXECUTABLE_INSTALL_RULE)
set(oneValueArgs DISABLE_TESTS RESULT_CACHE_FILE)
set(multiValueArgs DEPENDS DEPENDEES TEST_COMMAND_OPTIONS)
cmake_parse_arguments(_ "${options}" "${oneValueArgs}" "${multiValueArgs}" ${ARGN})
if (NOT DEFINED __DISABLE_TESTS)
set(__DISABLE_TESTS OFF)
endif()
if (__RESULT_CACHE_FILE)
add_custom_command(
TARGET ${TARGET}
POST_BUILD
COMMAND ${CMAKE_COMMAND}
ARGS -E copy ${__RESULT_CACHE_FILE} "$<TARGET_FILE_DIR:${TARGET}>"
)
endif()
if (NOT __DISABLE_EXECUTABLE_INSTALL_RULE AND CUTLASS_INSTALL_TESTS)
# file(RELATIVE_PATH CMAKE_CURRENT_BINARY_RELATIVE_DIR ${CMAKE_BINARY_DIR} ${CMAKE_CURRENT_BINARY_DIR})
install(
TARGETS ${TARGET}
RUNTIME DESTINATION ${CUTLASS_TEST_INSTALL_BINDIR}
)
if (__RESULT_CACHE_FILE)
install(
FILES ${__RESULT_CACHE_FILE}
DESTINATION ${CUTLASS_TEST_INSTALL_BINDIR}/
)
endif()
endif()
if (NOT __TEST_COMMAND_OPTIONS)
set(__TEST_COMMAND_OPTIONS " ")
endif()
list(LENGTH __TEST_COMMAND_OPTIONS CMD_COUNT)
set(CMD_IDX 0)
if (CMD_COUNT GREATER 1)
add_custom_target(${NAME} DEPENDS ${TARGET} ${__DEPENDS})
foreach(DEPENDEE ${__DEPENDEES})
add_dependencies(${DEPENDEE} ${NAME})
endforeach()
endif()
foreach(CMD_OPTIONS ${__TEST_COMMAND_OPTIONS})
if (CMD_COUNT GREATER 1)
set(TEST_NAME ${NAME}_${CMD_IDX})
else()
set(TEST_NAME ${NAME})
endif()
# The following rigmarole is needed to deal with spaces and possible quotes in
# command line arguments. The options are passed "by reference" as the actual
# variable names holding the real options. We then expand these in a way that
# preserves any quotes. Note, they have to be in this order for it to work for
# all the use cases below.
set(CMD_OPTIONS ${${CMD_OPTIONS}})
list(JOIN CMD_OPTIONS " " TEST_COMMAND_OPTIONS)
separate_arguments(CMD_OPTIONS)
add_custom_target(
${TEST_NAME}
COMMAND
${CUTLASS_TEST_EXECUTION_ENVIRONMENT} $<TARGET_FILE:${TARGET}> ${CMD_OPTIONS}
DEPENDS
${TARGET}
)
if (CMD_COUNT GREATER 1)
add_dependencies(${NAME} ${TEST_NAME})
endif()
foreach(DEPENDEE ${__DEPENDEES})
add_dependencies(${DEPENDEE} ${TEST_NAME})
endforeach()
add_test(
NAME c${TEST_NAME}
COMMAND ${CUTLASS_TEST_EXECUTION_ENVIRONMENT} $<TARGET_FILE:${TARGET}> ${CMD_OPTIONS}
)
set_tests_properties(c${TEST_NAME} PROPERTIES DISABLED ${__DISABLE_TESTS})
if (CUTLASS_INSTALL_TESTS)
# To run the tests from an install package with tests enabled, we need to generate test files
# that don't rely on the current directory structure in build.
set(TEST_NAME c${TEST_NAME})
set(TEST_EXE $<TARGET_FILE_NAME:${TARGET}>)
set(TEST_EXE_WORKING_DIRECTORY ./${CMAKE_INSTALL_BINDIR})
configure_file("${CUTLASS_CTEST_TEMPLATE_FILE}" "${CMAKE_PROJECT_DIR}${CMAKE_CURRENT_BINARY_DIR}/CTestTestfile.${TEST_NAME}.config.cmake" @ONLY)
file(GENERATE
OUTPUT "${CMAKE_PROJECT_DIR}${CMAKE_CURRENT_BINARY_DIR}/CTestTestfile.${TEST_NAME}.cmake"
INPUT "${CMAKE_PROJECT_DIR}${CMAKE_CURRENT_BINARY_DIR}/CTestTestfile.${TEST_NAME}.config.cmake"
)
install(
FILES "${CMAKE_PROJECT_DIR}${CMAKE_CURRENT_BINARY_DIR}/CTestTestfile.${TEST_NAME}.cmake"
DESTINATION ${CUTLASS_TEST_INSTALL_PREFIX}/ctest/
)
set(CUTLASS_CTEST_GENERATED_FILES ${CUTLASS_CTEST_GENERATED_FILES};ctest/CTestTestfile.${TEST_NAME}.cmake CACHE INTERNAL "")
endif()
math(EXPR CMD_IDX "${CMD_IDX} + 1")
endforeach()
endfunction()
if (CUTLASS_ENABLE_TOOLS)
add_subdirectory(tools)
if (CUTLASS_ENABLE_PROFILER)
add_dependencies(test_all test_profiler)
endif()
endif()
if (CUTLASS_ENABLE_EXAMPLES)
add_subdirectory(examples)
add_dependencies(test_all test_examples)
endif()
if (CUTLASS_ENABLE_TESTS)
add_subdirectory(test)
add_dependencies(test_all test_unit)
endif()
if (CUTLASS_INSTALL_TESTS)
file(MAKE_DIRECTORY "${CMAKE_BINARY_DIR}/cmake")
file(WRITE "${CMAKE_BINARY_DIR}/cmake/CTestTestfile.cmake" "# Generated File\n")
foreach(GENERATED_FILE ${CUTLASS_CTEST_GENERATED_FILES})
file(APPEND "${CMAKE_BINARY_DIR}/cmake/CTestTestfile.cmake" "include(${GENERATED_FILE})\n")
endforeach()
install(
FILES "${CMAKE_BINARY_DIR}/cmake/CTestTestfile.cmake"
DESTINATION "${CUTLASS_TEST_INSTALL_PREFIX}/"
)
endif()
#? install(
#? FILES ${CMAKE_BINARY_DIR}/CTestTestfile.cmake
#? DESTINATION ${CUTLASS_TEST_INSTALL_PREFIX}/
#? )
#?
#? install(
#? DIRECTORY
#? ${CMAKE_BINARY_DIR}/tools
#? ${CMAKE_BINARY_DIR}/test
#? DESTINATION ${CUTLASS_TEST_INSTALL_PREFIX}/
#? FILES_MATCHING PATTERN "CTestTestfile.cmake"
#? )
################################################################################
include(CMakePackageConfigHelpers)
write_basic_package_version_file(
${CMAKE_CURRENT_BINARY_DIR}/NvidiaCutlassConfigVersion.cmake
COMPATIBILITY AnyNewerVersion)
install(
FILES
${CMAKE_CURRENT_SOURCE_DIR}/cmake/NvidiaCutlassConfig.cmake
${CMAKE_CURRENT_BINARY_DIR}/NvidiaCutlassConfigVersion.cmake
DESTINATION ${CMAKE_INSTALL_LIBDIR}/cmake/NvidiaCutlass/
)
install(
EXPORT NvidiaCutlass
NAMESPACE nvidia::cutlass::
DESTINATION ${CMAKE_INSTALL_LIBDIR}/cmake/NvidiaCutlass/
FILE NvidiaCutlassTargets.cmake
)
################################################################################
include(${CMAKE_CURRENT_SOURCE_DIR}/cmake/NvidiaCutlassPackageConfig.cmake)

83
CONTRIBUTORS.md Normal file
View File

@ -0,0 +1,83 @@
![ALT](/media/images/gemm-hierarchy-with-epilogue-no-labels.png "CUTLASS")
[README](/README.md#documentation) > **Contributors**
# CUTLASS Developers and Contributors
This is the official list of CUTLASS developers and contributors.
## DEVELOPERS
Vijay Thakkar<br />
Pradeep Ramani<br />
Cris Cecka<br />
Aniket Shivam<br />
Jack Kosaian<br />
Mark Hoemmen<br />
Honghao Lu<br />
Ethan Yan<br />
Haicheng Wu<br />
Andrew Kerr<br />
Dustyn Blasig<br />
Fengqi Qiao<br />
Duane Merrill<br />
Yujia Zhai<br />
Shang Zhang<br />
Piotr Majcher<br />
Paul Springer<br />
Markus Hohnerbach<br />
Jin Wang<br />
Aditya Atluri<br />
## CuTe
Cris Cecka<br />
Vijay Thakkar<br />
## CUTLASS Product Manager
Matthew Nicely<br />
## Former CUTLASS Developers
Manish Gupta<br />
Naila Farooqui<br />
David Tanner<br />
Manikandan Ananth<br />
Zhaodong Chen<br />
Chinmay Talegaonkar<br />
## CONTRIBUTORS
Timothy Costa<br />
Julien Demouth<br />
Brian Fahs<br />
Michael Garland<br />
Michael Goldfarb<br />
Mostafa Hagog<br />
Fei Hu<br />
Alan Kaatz<br />
Tina Li<br />
Timmy Liu<br />
Wei Liu<br />
Duane Merrill<br />
Kevin Siu<br />
Markus Tavenrath<br />
John Tran<br />
Vicki Wang<br />
Junkai Wu<br />
Fung Xie<br />
Albert Xu<br />
Yang Xu<br />
Jack Yang<br />
Scott Yokim<br />
Xiuxia Zhang<br />
Nick Zhao<br />
## ACKNOWLEDGEMENTS
Girish Bharambe<br />
Luke Durant<br />
Carter Edwards<br />
Olivier Giroux<br />
Stephen Jones<br />
Rishkul Kulkarni<br />
Bryce Lelbach<br />
Joel McCormack<br />
Kyrylo Perelygin<br />
Sean Treichler<br />

362
CUDA.cmake Normal file
View File

@ -0,0 +1,362 @@
# Copyright (c) 2017 - 2023 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: BSD-3-Clause
#
# Redistribution and use in source and binary forms, with or without
# modification, are permitted provided that the following conditions are met:
#
# 1. Redistributions of source code must retain the above copyright notice, this
# list of conditions and the following disclaimer.
#
# 2. Redistributions in binary form must reproduce the above copyright notice,
# this list of conditions and the following disclaimer in the documentation
# and/or other materials provided with the distribution.
#
# 3. Neither the name of the copyright holder nor the names of its
# contributors may be used to endorse or promote products derived from
# this software without specific prior written permission.
#
# THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
# AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
# IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
# DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
# FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
# DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
# SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
# CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
# OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
# OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
if(CUDA_COMPILER MATCHES "[Cc]lang")
set(CUTLASS_NATIVE_CUDA_INIT ON)
elseif(CMAKE_VERSION VERSION_LESS 3.12.4)
set(CUTLASS_NATIVE_CUDA_INIT OFF)
else()
set(CUTLASS_NATIVE_CUDA_INIT ON)
endif()
set(CUTLASS_NATIVE_CUDA ${CUTLASS_NATIVE_CUDA_INIT} CACHE BOOL "Utilize the CMake native CUDA flow")
if(NOT DEFINED ENV{CUDACXX} AND NOT DEFINED ENV{CUDA_BIN_PATH} AND DEFINED ENV{CUDA_PATH})
# For backward compatibility, allow use of CUDA_PATH.
set(ENV{CUDACXX} $ENV{CUDA_PATH}/bin/nvcc)
endif()
if(CUTLASS_NATIVE_CUDA)
enable_language(CUDA)
if(NOT CUDA_VERSION)
set(CUDA_VERSION ${CMAKE_CUDA_COMPILER_VERSION})
endif()
if(NOT CUDA_TOOLKIT_ROOT_DIR)
get_filename_component(CUDA_TOOLKIT_ROOT_DIR "${CMAKE_CUDA_COMPILER}/../.." ABSOLUTE)
endif()
else()
find_package(CUDA REQUIRED)
# We workaround missing variables with the native flow by also finding the CUDA toolkit the old way.
if(NOT CMAKE_CUDA_COMPILER_VERSION)
set(CMAKE_CUDA_COMPILER_VERSION ${CUDA_VERSION})
endif()
endif()
if (CUDA_VERSION VERSION_LESS 9.2)
message(FATAL_ERROR "CUDA 9.2+ Required, Found ${CUDA_VERSION}.")
endif()
if(NOT CUTLASS_NATIVE_CUDA OR CUDA_COMPILER MATCHES "[Cc]lang")
set(CMAKE_CUDA_COMPILER ${CUDA_TOOLKIT_ROOT_DIR}/bin/nvcc)
message(STATUS "CUDA Compiler: ${CMAKE_CUDA_COMPILER}")
endif()
find_library(
CUDART_LIBRARY cudart
PATHS
${CUDA_TOOLKIT_ROOT_DIR}
PATH_SUFFIXES
lib/x64
lib64
lib
NO_DEFAULT_PATH
# We aren't going to search any system paths. We want to find the runtime
# in the CUDA toolkit we're building against.
)
if(NOT TARGET cudart AND CUDART_LIBRARY)
message(STATUS "CUDART: ${CUDART_LIBRARY}")
if(WIN32)
add_library(cudart STATIC IMPORTED GLOBAL)
# Even though we're linking against a .dll, in Windows you statically link against
# the .lib file found under lib/x64. The .dll will be loaded at runtime automatically
# from the PATH search.
else()
add_library(cudart SHARED IMPORTED GLOBAL)
endif()
add_library(nvidia::cudart ALIAS cudart)
set_property(
TARGET cudart
PROPERTY IMPORTED_LOCATION
${CUDART_LIBRARY}
)
elseif(TARGET cudart)
message(STATUS "CUDART: Already Found")
else()
message(STATUS "CUDART: Not Found")
endif()
find_library(
CUDA_DRIVER_LIBRARY cuda
PATHS
${CUDA_TOOLKIT_ROOT_DIR}
PATH_SUFFIXES
lib/x64
lib64
lib
lib64/stubs
lib/stubs
NO_DEFAULT_PATH
# We aren't going to search any system paths. We want to find the runtime
# in the CUDA toolkit we're building against.
)
if(NOT TARGET cuda_driver AND CUDA_DRIVER_LIBRARY)
message(STATUS "CUDA Driver: ${CUDA_DRIVER_LIBRARY}")
if(WIN32)
add_library(cuda_driver STATIC IMPORTED GLOBAL)
# Even though we're linking against a .dll, in Windows you statically link against
# the .lib file found under lib/x64. The .dll will be loaded at runtime automatically
# from the PATH search.
else()
add_library(cuda_driver SHARED IMPORTED GLOBAL)
endif()
add_library(nvidia::cuda_driver ALIAS cuda_driver)
set_property(
TARGET cuda_driver
PROPERTY IMPORTED_LOCATION
${CUDA_DRIVER_LIBRARY}
)
elseif(TARGET cuda_driver)
message(STATUS "CUDA Driver: Already Found")
else()
message(STATUS "CUDA Driver: Not Found")
endif()
find_library(
NVRTC_LIBRARY nvrtc
PATHS
${CUDA_TOOLKIT_ROOT_DIR}
PATH_SUFFIXES
lib/x64
lib64
lib
NO_DEFAULT_PATH
# We aren't going to search any system paths. We want to find the runtime
# in the CUDA toolkit we're building against.
)
if(NOT TARGET nvrtc AND NVRTC_LIBRARY)
message(STATUS "NVRTC: ${NVRTC_LIBRARY}")
if(WIN32)
add_library(nvrtc STATIC IMPORTED GLOBAL)
# Even though we're linking against a .dll, in Windows you statically link against
# the .lib file found under lib/x64. The .dll will be loaded at runtime automatically
# from the PATH search.
else()
add_library(nvrtc SHARED IMPORTED GLOBAL)
endif()
add_library(nvidia::nvrtc ALIAS nvrtc)
set_property(
TARGET nvrtc
PROPERTY IMPORTED_LOCATION
${NVRTC_LIBRARY}
)
elseif(TARGET nvrtc)
message(STATUS "NVRTC: Already Found")
else()
message(STATUS "NVRTC: Not Found")
endif()
include_directories(SYSTEM ${CUDA_INCLUDE_DIRS})
# Some platforms (e.g. Visual Studio) don't add the CUDA include directories to the system include
# paths by default, so we add it explicitly here.
function(cutlass_correct_source_file_language_property)
if(CUDA_COMPILER MATCHES "[Cc]lang")
foreach(File ${ARGN})
if(File MATCHES ".*\.cu$")
set_source_files_properties(${File} PROPERTIES LANGUAGE CXX)
endif()
endforeach()
endif()
endfunction()
if (MSVC OR CUTLASS_LIBRARY_KERNELS MATCHES "all")
set(CUTLASS_UNITY_BUILD_ENABLED_INIT ON)
else()
set(CUTLASS_UNITY_BUILD_ENABLED_INIT OFF)
endif()
set(CUTLASS_UNITY_BUILD_ENABLED ${CUTLASS_UNITY_BUILD_ENABLED_INIT} CACHE BOOL "Enable combined source compilation")
set(CUTLASS_UNITY_BUILD_BATCH_SIZE 16 CACHE STRING "Batch size for unified source files")
function(cutlass_unify_source_files TARGET_ARGS_VAR)
set(options)
set(oneValueArgs BATCH_SOURCES BATCH_SIZE)
set(multiValueArgs)
cmake_parse_arguments(_ "${options}" "${oneValueArgs}" "${multiValueArgs}" ${ARGN})
if (NOT DEFINED TARGET_ARGS_VAR)
message(FATAL_ERROR "TARGET_ARGS_VAR parameter is required")
endif()
if (__BATCH_SOURCES AND NOT DEFINED __BATCH_SIZE)
set(__BATCH_SIZE ${CUTLASS_UNITY_BUILD_BATCH_SIZE})
endif()
if (CUTLASS_UNITY_BUILD_ENABLED AND DEFINED __BATCH_SIZE AND __BATCH_SIZE GREATER 1)
set(CUDA_FILE_ARGS)
set(TARGET_SOURCE_ARGS)
foreach(ARG ${__UNPARSED_ARGUMENTS})
if(${ARG} MATCHES ".*\.cu$")
list(APPEND CUDA_FILE_ARGS ${ARG})
else()
list(APPEND TARGET_SOURCE_ARGS ${ARG})
endif()
endforeach()
list(LENGTH CUDA_FILE_ARGS NUM_CUDA_FILE_ARGS)
while(NUM_CUDA_FILE_ARGS GREATER 0)
list(SUBLIST CUDA_FILE_ARGS 0 ${__BATCH_SIZE} CUDA_FILE_BATCH)
string(SHA256 CUDA_FILE_BATCH_HASH "${CUDA_FILE_BATCH}")
string(SUBSTRING ${CUDA_FILE_BATCH_HASH} 0 12 CUDA_FILE_BATCH_HASH)
set(BATCH_FILE ${CMAKE_CURRENT_BINARY_DIR}/${NAME}.unity.${CUDA_FILE_BATCH_HASH}.cu)
message(STATUS "Generating ${BATCH_FILE}")
file(WRITE ${BATCH_FILE} "// Unity File - Auto Generated!\n")
foreach(CUDA_FILE ${CUDA_FILE_BATCH})
get_filename_component(CUDA_FILE_ABS_PATH ${CUDA_FILE} ABSOLUTE)
file(APPEND ${BATCH_FILE} "#include \"${CUDA_FILE_ABS_PATH}\"\n")
endforeach()
list(APPEND TARGET_SOURCE_ARGS ${BATCH_FILE})
if (NUM_CUDA_FILE_ARGS LESS_EQUAL __BATCH_SIZE)
break()
endif()
list(SUBLIST CUDA_FILE_ARGS ${__BATCH_SIZE} -1 CUDA_FILE_ARGS)
list(LENGTH CUDA_FILE_ARGS NUM_CUDA_FILE_ARGS)
endwhile()
else()
set(TARGET_SOURCE_ARGS ${__UNPARSED_ARGUMENTS})
endif()
set(${TARGET_ARGS_VAR} ${TARGET_SOURCE_ARGS} PARENT_SCOPE)
endfunction()
function(cutlass_add_library NAME)
set(options SKIP_GENCODE_FLAGS)
set(oneValueArgs EXPORT_NAME)
set(multiValueArgs)
cmake_parse_arguments(_ "${options}" "${oneValueArgs}" "${multiValueArgs}" ${ARGN})
cutlass_unify_source_files(TARGET_SOURCE_ARGS ${__UNPARSED_ARGUMENTS})
if(CUTLASS_NATIVE_CUDA OR CUDA_COMPILER MATCHES "clang")
cutlass_correct_source_file_language_property(${TARGET_SOURCE_ARGS})
add_library(${NAME} ${TARGET_SOURCE_ARGS})
else()
set(CUDA_LINK_LIBRARIES_KEYWORD PRIVATE)
cuda_add_library(${NAME} ${TARGET_SOURCE_ARGS})
endif()
cutlass_apply_standard_compile_options(${NAME})
if (NOT __SKIP_GENCODE_FLAGS)
cutlass_apply_cuda_gencode_flags(${NAME})
endif()
target_compile_features(
${NAME}
INTERFACE
cxx_std_11
)
if(__EXPORT_NAME)
add_library(nvidia::cutlass::${__EXPORT_NAME} ALIAS ${NAME})
set_target_properties(${NAME} PROPERTIES EXPORT_NAME ${__EXPORT_NAME})
endif()
endfunction()
function(cutlass_add_executable NAME)
set(options)
set(oneValueArgs)
set(multiValueArgs)
cmake_parse_arguments(_ "${options}" "${oneValueArgs}" "${multiValueArgs}" ${ARGN})
cutlass_unify_source_files(TARGET_SOURCE_ARGS ${__UNPARSED_ARGUMENTS})
if(CUTLASS_NATIVE_CUDA OR CUDA_COMPILER MATCHES "clang")
cutlass_correct_source_file_language_property(${TARGET_SOURCE_ARGS})
add_executable(${NAME} ${TARGET_SOURCE_ARGS})
else()
set(CUDA_LINK_LIBRARIES_KEYWORD PRIVATE)
cuda_add_executable(${NAME} ${TARGET_SOURCE_ARGS})
endif()
cutlass_apply_standard_compile_options(${NAME})
cutlass_apply_cuda_gencode_flags(${NAME})
target_compile_features(
${NAME}
INTERFACE
cxx_std_11
)
endfunction()
function(cutlass_target_sources NAME)
set(options)
set(oneValueArgs)
set(multiValueArgs)
cmake_parse_arguments(_ "${options}" "${oneValueArgs}" "${multiValueArgs}" ${ARGN})
cutlass_unify_source_files(TARGET_SOURCE_ARGS ${__UNPARSED_ARGUMENTS})
cutlass_correct_source_file_language_property(${TARGET_SOURCE_ARGS})
target_sources(${NAME} ${TARGET_SOURCE_ARGS})
endfunction()

View File

@ -32,7 +32,7 @@ DOXYFILE_ENCODING = UTF-8
# title of most generated pages and in a few other places.
# The default value is: My Project.
PROJECT_NAME = "Cutlass"
PROJECT_NAME = "CUTLASS"
# The PROJECT_NUMBER tag can be used to enter a project or revision number. This
# could be handy for archiving the generated documentation or if some version
@ -51,7 +51,7 @@ PROJECT_BRIEF = "CUDA Templates for Linear Algebra Subroutines and Solv
# and the maximum width should not exceed 200 pixels. Doxygen will copy the logo
# to the output directory.
PROJECT_LOGO =
PROJECT_LOGO = media/images/cutlass-logo-small.png
# The OUTPUT_DIRECTORY tag is used to specify the (relative or absolute) path
# into which the generated documentation will be written. If a relative path is
@ -206,7 +206,7 @@ SEPARATE_MEMBER_PAGES = NO
# uses this value to replace tabs by spaces in code fragments.
# Minimum value: 1, maximum value: 16, default value: 4.
TAB_SIZE = 4
TAB_SIZE = 2
# This tag can be used to specify a number of aliases that act as commands in
# the documentation. An alias has the form:
@ -218,7 +218,8 @@ TAB_SIZE = 4
# "Side Effects:". You can put \n's in the value part of an alias to insert
# newlines.
ALIASES =
#ALIASES += "concept{1}=@ingroup \1\n@par Implemented concepts:\n@ref \1"
ALIASES += "concept{1}=@ingroup \1"
# This tag can be used to specify a number of word-keyword mappings (TCL only).
# A mapping has the form "name=value". For example adding "class=itcl::class"
@ -296,7 +297,7 @@ AUTOLINK_SUPPORT = YES
# diagrams that involve STL classes more complete and accurate.
# The default value is: NO.
BUILTIN_STL_SUPPORT = NO
BUILTIN_STL_SUPPORT = YES
# If you use Microsoft's C++/CLI language, you should set this option to YES to
# enable parsing support.
@ -396,7 +397,7 @@ LOOKUP_CACHE_SIZE = 0
# normally produced when WARNINGS is set to YES.
# The default value is: NO.
EXTRACT_ALL = NO
EXTRACT_ALL = YES
# If the EXTRACT_PRIVATE tag is set to YES all private members of a class will
# be included in the documentation.
@ -733,7 +734,9 @@ WARN_LOGFILE =
# spaces.
# Note: If this tag is empty the current directory is searched.
INPUT = cutlass cutlass/gemm cutlass/util
INPUT = include/cutlass tools/util/include/cutlass/ tools/library/include/cutlass/
INPUT += media/docs/doxygen_mainpage.md
# This tag can be used to specify the character encoding of the source files
# that doxygen parses. Internally doxygen uses the UTF-8 encoding. Doxygen uses
@ -759,7 +762,7 @@ FILE_PATTERNS =
# be searched for input files as well.
# The default value is: NO.
RECURSIVE = NO
RECURSIVE = YES
# The EXCLUDE tag can be used to specify files and/or directories that should be
# excluded from the INPUT source files. This way you can easily exclude a
@ -869,7 +872,7 @@ FILTER_SOURCE_PATTERNS =
# (index.html). This can be useful if you have a project on for instance GitHub
# and want to reuse the introduction page also for the doxygen output.
USE_MDFILE_AS_MAINPAGE =
USE_MDFILE_AS_MAINPAGE = media/docs/doxygen_mainpage.md
#---------------------------------------------------------------------------
# Configuration options related to source browsing
@ -998,7 +1001,7 @@ GENERATE_HTML = YES
# The default directory is: html.
# This tag requires that the tag GENERATE_HTML is set to YES.
HTML_OUTPUT = generated-html
HTML_OUTPUT =
# The HTML_FILE_EXTENSION tag can be used to specify the file extension for each
# generated HTML page (for example: .htm, .php, .asp).
@ -1079,7 +1082,7 @@ HTML_EXTRA_FILES =
# Minimum value: 0, maximum value: 359, default value: 220.
# This tag requires that the tag GENERATE_HTML is set to YES.
HTML_COLORSTYLE_HUE = 82
HTML_COLORSTYLE_HUE = 100
# The HTML_COLORSTYLE_SAT tag controls the purity (or saturation) of the colors
# in the HTML output. For a value of 0 the output will use grayscales only. A
@ -1087,7 +1090,7 @@ HTML_COLORSTYLE_HUE = 82
# Minimum value: 0, maximum value: 255, default value: 100.
# This tag requires that the tag GENERATE_HTML is set to YES.
HTML_COLORSTYLE_SAT = 100
HTML_COLORSTYLE_SAT = 50
# The HTML_COLORSTYLE_GAMMA tag controls the gamma correction applied to the
# luminance component of the colors in the HTML output. Values below 100
@ -1106,7 +1109,7 @@ HTML_COLORSTYLE_GAMMA = 80
# The default value is: YES.
# This tag requires that the tag GENERATE_HTML is set to YES.
HTML_TIMESTAMP = YES
HTML_TIMESTAMP = NO
# If the HTML_DYNAMIC_SECTIONS tag is set to YES then the generated HTML
# documentation will contain sections that can be hidden and shown after the
@ -2032,7 +2035,7 @@ HIDE_UNDOC_RELATIONS = YES
# set to NO
# The default value is: NO.
HAVE_DOT = NO
HAVE_DOT = $(HAVE_DOT)
# The DOT_NUM_THREADS specifies the number of dot invocations doxygen is allowed
# to run in parallel. When set to 0 doxygen will base this on the number of
@ -2204,7 +2207,7 @@ INTERACTIVE_SVG = NO
# found. If left blank, it is assumed the dot tool can be found in the path.
# This tag requires that the tag HAVE_DOT is set to YES.
DOT_PATH =
DOT_PATH = $(DOT_PATH)
# The DOTFILE_DIRS tag can be used to specify one or more directories that
# contain dot files that are included in the documentation (see the \dotfile

27
LICENSE.txt Normal file
View File

@ -0,0 +1,27 @@
Copyright (c) 2017 - 2023 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
SPDX-License-Identifier: BSD-3-Clause
Redistribution and use in source and binary forms, with or without
modification, are permitted provided that the following conditions are met:
1. Redistributions of source code must retain the above copyright notice, this
list of conditions and the following disclaimer.
2. Redistributions in binary form must reproduce the above copyright notice,
this list of conditions and the following disclaimer in the documentation
and/or other materials provided with the distribution.
3. Neither the name of the copyright holder nor the names of its
contributors may be used to endorse or promote products derived from
this software without specific prior written permission.
THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.

30
PUBLICATIONS.md Normal file
View File

@ -0,0 +1,30 @@
# Publications Using Cutlass
## 2023
- ["Stream-K: Work-centric Parallel Decomposition for Dense Matrix-Matrix Multiplication on the GPU"](https://arxiv.org/abs/2301.03598). Muhammad Osama, Duane Merrill, Cris Cecka, Michael Garland, John D. Owens. _arXiv_, January 2023.
## 2022
- ["GPU Load Balancing"](https://arxiv.org/abs/2212.08964). Muhammad Osama. _Doctoral dissertation, University of California, Davis_, December 2022.
- ["Bolt: Bridging the Gap between Auto-tuners and Hardware-native Performance"](https://arxiv.org/abs/2110.15238). Jiarong Xing, Leyuan Wang, Shang Zhang, Jack Chen, Ang Chen, Yibo Zhu. _Proceedings of the 5th MLSys Conference_, August 2022.
- ["Recovering single precision accuracy from Tensor Cores while surpassing the FP32 theoretical peak performance"](https://arxiv.org/abs/2203.03341). Hiroyuki Ootomo, Rio Yokota. _International Journal of High Performance Computing_, March 2022.
- ["Breaking the Computation and Communication Abstraction Barrier in Distributed Machine Learning Workloads"](https://arxiv.org/abs/2105.05720). Abhinav Jangda, Jun Huang, Guodong Liu, Amir Hossein Nodehi Sabet, Saeed Maleki, Youshan Miao, Madanlal Musuvathi, Todd Mytkowicz, Olli Sarikivi. _Proceedings of the 27th ACM International Conference on Architectural Support for Programming Languages and Operating Systems_, February 2022.
## 2021
- ["Arithmetic-intensity-guided fault tolerance for neural network inference on GPUs"](https://dl.acm.org/doi/abs/10.1145/3458817.3476184). Jack Kosaian, K. V. Rashmi. _Proceedings of the International Conference for High Performance Computing, Networking, Storage and Analysis_, November 2021.
- ["Real-time Neural Radiance Caching for Path Tracing"](https://d1qx31qr3h6wln.cloudfront.net/publications/paper_4.pdf). Thomas Muller, Fabrice Rousselle, Jan Novak, Alex Keller. _ACM Trans. Graph._, August 2021.
## 2020
- ["Scalable Knowledge Graph Analytics at 136 Petaflop/s"](https://www.computer.org/csdl/proceedings-article/sc/2020/999800a061/1oeORDgCM0g). Ramakrishnan Kannan, Piyush Sao, Hao Lu, Drahomira Herrmannova, Vijay Thakkar, Robert Patton, Richard Vuduc, Thomas Potok. _Proceedings of the International Conference for High Performance Computing, Networking, Storage and Analysis_, November 2020.
- ["Accelerating Sparse DNN Models without Hardware-Support via Tile-Wise Sparsity
"](https://arxiv.org/abs/2008.13006). Cong Guo, Bo Yang Hsueh, Jingwen Leng, Yuxian Qiu, Yue Guan, Zehuan Wang, Xiaoying Jia, Xipeng Li, Minyi Guo, Yuhao Zhu. _Proceedings of the International Conference for High Performance Computing, Networking, Storage and Analysis_, November 2020.
- ["Strassen's Algorithm Reloaded on GPUs"](https://dl.acm.org/doi/10.1145/3372419). Jianyu Huang, Chenhan D. Yu, Robert A. van de Geijn. _ACM Transactions on Mathematical Software_, March 2020.

584
README.md
View File

@ -1,75 +1,561 @@
# Introduction
![ALT](/media/images/gemm-hierarchy-with-epilogue-no-labels.png "Complete CUDA GEMM decomposition")
# CUTLASS 3.0
_CUTLASS 3.0 - January 2023_
CUTLASS is a collection of CUDA C++ template abstractions for implementing
high-performance matrix-matrix multiplication (GEMM) and related computations at all levels
and scales within CUDA. It incorporates strategies for hierarchical decomposition and
data movement similar to those used to implement cuBLAS and cuDNN. CUTLASS decomposes
these "moving parts" into reusable, modular software components abstracted by C++ template
classes. Primitives for different levels of a conceptual parallelization hierarchy
can be specialized and tuned via custom tiling sizes, data types,
and other algorithmic policy. The resulting flexibility simplifies their use
as building blocks within custom kernels and applications.
To support a wide variety of applications, CUTLASS provides extensive support for
mixed-precision computations, providing specialized data-movement and
multiply-accumulate abstractions for half-precision floating
point (FP16), BFloat16 (BF16), Tensor Float 32 (TF32),
single-precision floating point (FP32),
[FP32 emulation via tensor core instruction](/examples/27_ampere_3xtf32_fast_accurate_tensorop_gemm),
double-precision floating
point (FP64) types, integer data types (4b and 8b), and binary data types (1b).
CUTLASS demonstrates warp-synchronous matrix multiply operations
targeting the programmable, high-throughput _Tensor Cores_ implemented by
NVIDIA's Volta, Turing, Ampere, and Hopper architectures.
See the [Quick Start Guide](/media/docs/quickstart.md) to get started quickly.
See the [functionality listing](/media/docs/functionality.md) for the list of operations
supported at each level of the execution model hierarchy.
CUTLASS 3.0 introduces a new core library, CuTe, to describe and manipulate tensors of threads and data.
CuTe is a collection of C++ CUDA template abstractions for defining and operating on hierarchically multidimensional layouts of threads and data. CuTe provides `Layout` and `Tensor` objects that compactly package the type, shape, memory space, and layout of data, while performing the complicated indexing for the user. This lets programmers focus on the logical descriptions of their algorithms while CuTe does the mechanical bookkeeping for them. With these tools, we can quickly design, implement, and modify all dense linear algebra operations.
The core abstractions of CuTe are hierarchically multidimensional layouts which can be composed with data arrays to represent tensors. The representation of layouts is powerful enough to represent nearly everything we need to implement efficient dense linear algebra. Layouts can also be combined and manipulated via functional composition, on which we build a large set of common operations such as tiling and partitioning.
CUTLASS 3.0 adopts CuTe throughout the GEMM hierarchy in its templates. This greatly simplifies the design
and improves code composability and readability. More documentation specific to CuTe can be found in its [dedicated documentation directory](/media/docs/cute/00_quickstart.md).
In addition to GEMMs, CUTLASS implements high-performance convolution via the implicit GEMM algorithm. Implicit GEMM is the formulation of a convolution operation as a GEMM thereby taking advantage of CUTLASS's modular GEMM pipeline. This allows CUTLASS to build convolutions by reusing highly-optimized GEMM components.
# What's New in CUTLASS 3.0
CUTLASS 3.0, as the next major version of the CUTLASS API, brings with it CuTe, a new programming model and backend designed for massively parallel heterogenous agents. Using CuTe, CUTLASS 3.0 provides implementations of GEMM kernels for the NVIDIA Hopper architecture.
- [CuTe-based layouts and layout algebra](/media/docs/cute/00_quickstart.md)
- [A new GEMM template API](/media/docs/gemm_api_3x.md) that eschews the architecture-centric hierarchy of 2.x in favour of a new conceptual framing. Read more in the [3.0 design documentation](/media/docs/cutlass_3x_design.md).
- Support for 4th generation Hopper Tensor Core instructions (WGMMA) through CuTe.
- Support for Hopper asynchronous Tensor Memory Accelerator (TMA) instructions and associated transaction barriers through CuTe.
- New warp-specialized GEMM kernels targeting Hopper TMA + WGMMA for speed-of-light GEMMs.
- New warp-specialized persistent GEMM kernels targeting Hopper TMA + WGMMA.
- Support for CUDA Threadblock Clusters and programmatic TMA multicast for greater execution and data locality.
- A new way to instantiate default GEMM kernels using `CollectiveBuilder`s that supersede the 2.x `DefaultXConfiguration` types in favour a metaprogramming based kernel generator functionality. See [example 49](/examples/49_hopper_gemm_schedules_with_collective_builder/49_hopper_gemm_schedules_with_collective_builder.cu).
- Extensions to the CUTLASS library and profiler to support CUTLASS 3.0 Hopper kernels, and a new format
for kernel procedural names.
- *Announcement*: CUTLASS plans to rename the GitHub branch `master` to `main` with a future release.
## New architecture, compiler, and CUDA Toolkit requirements
Minimum requirements:
- Architecture: Volta
- Compiler: Must support at least C++17
- CUDA Toolkit version: 11.4
CUTLASS 3.0 *removes support* for the following:
- Maxwell and Pascal GPU architectures
- Ubuntu 16.04
- CUDA 10.2
- C++ language versions less than 17.
**See the [CHANGELOG](CHANGELOG.md) for a detailed listing of releases and updates.**
# Performance
<p align="center"><img src=media/images/cutlass-3.0-gemm-peak-performance.png></p>
CUTLASS primitives are very efficient. When used to construct device-wide GEMM kernels,
they exhibit peak performance comparable to cuBLAS for scalar GEMM
computations. The above figure shows CUTLASS performance relative to cuBLAS
for large matrix dimensions on an [NVIDIA H100](https://www.nvidia.com/en-us/data-center/h100/) (NVIDIA Hopper architecture),
an [NVIDIA L40](https://www.nvidia.com/en-us/data-center/l40/) (NVIDIA Ada architecture),
an [NVIDIA A100](https://www.nvidia.com/en-us/data-center/a100/) (NVIDIA Ampere architecture),
and an [NVIDIA A40](https://www.nvidia.com/en-us/data-center/a40/) (NVIDIA Ampere architecture).
CUTLASS 3.0 was compiled with the [CUDA 12.0 Toolkit](https://developer.nvidia.com/cuda-downloads).
Tensor Core operations are implemented using CUDA's
[mma instruction](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#warp-level-matrix-instructions-mma).
<p align="center"><img src=media/images/cutlass-2.9-implicit-gemm-performance.png></p>
When using CUTLASS building blocks to construct device-wide implicit gemm (Fprop, Dgrad, and Wgrad)
kernels, CUTLASS performance is also comparable to cuDNN when running Resnet-50 layers on an [NVIDIA A100](https://www.nvidia.com/en-us/data-center/a100/)
as shown in the above figure. Tensor Core operations are still implemented using CUDA's
[mma instruction](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#warp-level-matrix-instructions-mma).
# Compatibility
CUTLASS requires a C++17 host compiler and
performs best when built with the [**CUDA 12.0 Toolkit**](https://developer.nvidia.com/cuda-toolkit).
It is also compatible with CUDA 11.4, CUDA 11.5, CUDA 11.6, CUDA 11.7, and CUDA 11.8.
## Operating Systems
We have tested the following environments.
|**Operating System** | **Compiler** |
|-----------------|----------|
| Ubuntu 18.04 | GCC 7.5.0 |
| Ubuntu 20.04 | GCC 10.3.0 |
| Ubuntu 22.04 | GCC 11.2.0 |
Note: We plan to add Windows (MSVC) & Clang compiler support soon.
## Hardware
CUTLASS runs successfully on the following NVIDIA GPUs, and it is expected to be efficient on Volta, Turing, Ampere, Ada, and Hopper architecture based NVIDIA GPUs.
|**GPU**|**CUDA Compute Capability**|**Minimum CUDA Toolkit Required by CUTLASS-3**|
|---|---|---|
|NVIDIA V100 Tensor Core GPU |7.0|11.4|
|NVIDIA TitanV |7.0|11.4|
|NVIDIA GeForce RTX 2080 TI, 2080, 2070 |7.5|11.4|
|NVIDIA T4 |7.5|11.4|
|NVIDIA A100 Tensor Core GPU |8.0|11.4|
|NVIDIA A10 |8.6|11.4|
|NVIDIA GeForce RTX 3090 |8.6|11.4|
|NVIDIA GeForce RTX 4090 |8.9|11.8|
|NVIDIA L40 |8.9|11.8|
|NVIDIA H100 Tensor Core GPU |9.0|11.8|
## Target Architecture
In general, PTX code generated for one target architecture can be run on future architectures (i.e., it is forward compatible). However, CUDA 12.0 introduces the concept of "architecture-accelerated features" whose PTX does not have forward compatibility guarantees. Several Hopper PTX instructions fall under this category of architecture-accelerated features, and thus require a `sm_90a` target architecture (note the "a" appended). For more details on this and other architecture-accelerated instructions, please refer to the [CUDA Documentation](https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#feature-availability).
The target architecture information is passed on to CUTLASS via the cmake flag `CUTLASS_NVCC_ARCHS`. In order to maximize performance on Hopper GH100, users are required to build CUTLASS with `90a` as the target architecture. If a user accidentally builds a kernel which uses SM90a features (e.g. Hopper Tensor Core Instructions), using the SM90 target (note the lack of "a"), with either CTK 12.0 or 11.8, the kernel is expected to fail with a runtime error.
```
cmake .. -DCUTLASS_NVCC_ARCHS="90a"
```
Please refer to the [functionality documentation](media/docs/functionality.md) for details on which kernels require which target architectures.
# Documentation
CUTLASS is described in the following documents and the accompanying
[Doxygen documentation](https://nvidia.github.io/cutlass).
- [Quick Start Guide](/media/docs/quickstart.md) - build and run CUTLASS
- [Functionality](/media/docs/functionality.md) - summarizes functionality available in CUTLASS
- [Efficient GEMM in CUDA](media/docs/efficient_gemm.md) - describes how GEMM kernels may be implemented efficiently in CUDA
- [CUTLASS 3.x Design](media/docs/cutlass_3x_design.md) - describes the CUTLASS 3.x design, its benefits, and how CuTe enables us to write much more composable components
- [GEMM API 3.x](media/docs/gemm_api_3x.md) - describes the CUTLASS 3.x GEMM model and C++ template concepts
- [GEMM API 2.x](media/docs/gemm_api.md) - describes the CUTLASS 2.x GEMM model and C++ template concepts
- [Implicit GEMM Convolution](media/docs/implicit_gemm_convolution.md) - describes 2-D and 3-D convolution in CUTLASS
- [Code Organization](media/docs/code_organization.md) - describes the organization and contents of the CUTLASS project
- [Terminology](media/docs/terminology.md) - describes terms used in the code
- [Programming Guidelines](media/docs/programming_guidelines.md) - guidelines for writing efficient modern CUDA C++
- [Fundamental types](media/docs/fundamental_types.md) - describes basic C++ classes used in CUTLASS to represent numeric quantities and arrays
- [Layouts](media/docs/layout.md) - describes layouts of matrices and tensors in memory
- [Tile Iterators](media/docs/tile_iterator_concept.md) - describes C++ concepts for iterating over tiles of matrices in memory
- [CUTLASS Profiler](media/docs/profiler.md) - command-line driven profiling application
- [CUTLASS Utilities](media/docs/utilities.md) - additional templates used to facilate rapid development
# Resources
We have also described the structure of an efficient GEMM in our talk at the
[GPU Technology Conference 2018](http://on-demand.gputechconf.com/gtc/2018/presentation/s8854-cutlass-software-primitives-for-dense-linear-algebra-at-all-levels-and-scales-within-cuda.pdf).
- [CUTLASS: Software Primitives for Dense Linear Algebra at All Levels and Scales within CUDA](https://www.nvidia.com/en-us/on-demand/session/gtcsiliconvalley2018-s8854/)
- [Developing CUDA Kernels to Push Tensor Cores to the Absolute Limit on NVIDIA A100](https://www.nvidia.com/en-us/on-demand/session/gtcsj20-s21745/)
- [Accelerating Convolution with Tensor Cores in CUTLASS](https://www.nvidia.com/en-us/on-demand/session/gtcspring21-s31883/)
- [Accelerating Backward Data Gradient by Increasing Tensor Core Utilization in CUTLASS](https://www.nvidia.com/en-us/on-demand/session/gtcspring22-s41996/)
- [CUTLASS: Python API, Enhancements, and NVIDIA Hopper](https://www.nvidia.com/en-us/on-demand/session/gtcfall22-a41131/)
# Building CUTLASS
CUTLASS is a header-only template library and does not need to be built to be used by other
projects. Client applications should target CUTLASS's `include/` directory in their include
paths.
CUTLASS unit tests, examples, and utilities can be build with CMake starting version 3.12.
Make sure the `CUDACXX` environment variable points to NVCC in the CUDA Toolkit installed
on your system.
```bash
$ export CUDACXX=${CUDA_INSTALL_PATH}/bin/nvcc
```
Create a build directory within the CUTLASS project, then run CMake. By default CUTLASS will build kernels
for CUDA architecture versions 5.0, 6.0, 6.1, 7.0, 7.5, 8.0, 8.6, 8.9, and 9.0.
To reduce compile time you can specify
the architectures to build CUTLASS for by changing the CMake configuration setting
`CUTLASS_NVCC_ARCHS`.
```bash
$ mkdir build && cd build
$ cmake .. -DCUTLASS_NVCC_ARCHS=80 # compiles for NVIDIA's Ampere Architecture
```
From the `build/` directory, compile and run the CUTLASS unit tests by building the target `test_unit` with make.
The unit tests are organized as several binaries mirroring the top-level namespaces of CUTLASS,
and they may be executed in parallel via make's `-j` command line argument.
```bash
$ make test_unit -j
...
...
...
[----------] Global test environment tear-down
[==========] 946 tests from 57 test cases ran. (10812 ms total)
[ PASSED ] 946 tests.
```
All tests should pass on supported platforms, though the exact number of tests may vary over time.
CUTLASS is a CUDA C++ template library for implementing matrix-multiply
procedures that may be instantiated in CUDA device kernels. CUTLASS applies
object-oriented and generic programming techniques to maximize flexibility of
the resulting code and facilitate composition with caller-supplied functionality.
CUDA C++ templates are used to specify policy decisions such as block sizes,
data types of input and accumulator operands, and element-wise operations applied
to the results of matrix multiply.
# Project Structure
CUTLASS is arranged as a header-only library with several example test programs
that demonstrate instantiating a GEMM task within a CUDA kernel. Comments inline
with the source explain the individual components.
CUTLASS is arranged as a header-only library along with Utilities, Tools, Examples, and unit tests.
[Doxygen documentation](https://nvidia.github.io/cutlass) provides a complete list of files, classes,
and template concepts defined in the CUTLASS project.
The repository is organized in the following arrangement.
A detailed explanation of the source code organization may be found in the
[CUTLASS documentation](media/docs/code_organization.md), but several main components are summarized below.
cutlass/ Root of header-only source library for matrix multiply
gemm/ Implementation of GEMM __device__ code and supporting components
util/ Utility components for CUDA device-side CUDA development
## CUTLASS Template Library
A test program is provided to illustrate the use of CUTLASS. This is implemented
in the following directory.
```
include/ # client applications should target this directory in their build's include paths
cutlass_test Root of test programs depicting CUTLASS kernels
util/ Utilities
gemm.cu Simple example calling CUTLASS and CUBLAS GEMM kernels
Makefile Build script for test programs
cutlass/ # CUDA Templates for Linear Algebra Subroutines and Solvers - headers only
arch/ # direct exposure of architecture features (including instruction-level GEMMs)
conv/ # code specialized for convolution
epilogue/ # code specialized for the epilogue of gemm/convolution
gemm/ # code specialized for general matrix product computations
layout/ # layout definitions for matrices, tensors, and other mathematical objects in memory
platform/ # CUDA-capable Standard Library components
reduction/ # bandwidth-limited reduction kernels that do not fit the "gemm" model
thread/ # simt code that can be performed within a CUDA thread
transform/ # code specialized for layout, type, and domain transformations
* # core vocabulary types, containers, and basic numeric operations
cute/ # CuTe Layout, layout algebra, MMA/Copy atoms, tiled MMA/Copy
algorithm/ # Definitions of core operations such as copy, gemm, and operations on cute::tuples
arch/ # Bare bones PTX wrapper structs for copy and math instructions
atom/ # Meta-information either link to or built from arch/ operators
mma_atom.hpp # cute::Mma_Atom and cute::TiledMma
copy_atom.hpp # cute::Copy_Atom and cute::TiledCopy
*sm*.hpp # Arch specific meta-information for copy and math operations
* # Core library types such as Shape, Stride, Layout, Tensor, and associated operations
```
### CUTLASS SDK Examples
[CUTLASS SDK examples](/examples) apply CUTLASS templates to implement basic computations.
### Tools
```
tools/
library/ # CUTLASS Instance Library - contains instantiations of all supported CUTLASS templates
include/
cutlass/
library/
profiler/ # CUTLASS Profiler - command-line utility for executing operations in the
# CUTLASS Library
util/ # CUTLASS Utilities - contains numerous helper classes for
include/ # manging tensors in device memory, reference
cutlass/ # implementations for GEMM, random initialization
util/ # of tensors, and I/O.
```
### Test
The `test/unit/` directory consist of unit tests implemented with Google Test that demonstrate
basic usage of Core API components and complete tests of the CUTLASS GEMM computations.
Instructions for building and running the Unit tests are described in the [Quickstart guide](media/docs/quickstart.md).
# Performance Profiling
The `tools/profiler/` directory contains a command-line utility for launching each of the GEMM kernels.
It can be built as follows:
```bash
$ make cutlass_profiler -j16
```
## Building all GEMM and Convolution kernels (_long_ build times)
By default, only one tile size is instantiated for each data type, math instruction, and layout.
To instantiate all, set the following environment variable when running CMake from an empty `build/` directory.
Beware, this results in *tens of thousands* of kernels and long build times.
This would also result in a large binary size and on some platforms linker to fail on building the library.
Therefore, it's highly recommended to generate only a subset of kernels as demonstrated in the sub-section below.
```bash
$ cmake .. -DCUTLASS_NVCC_ARCHS=90a -DCUTLASS_LIBRARY_KERNELS=all
...
$ make cutlass_profiler -j16
```
## Building a subset of GEMM and Convolution kernels (_reduced_ build times)
To compile strictly one kernel or a small set of kernels, a comma-delimited list of kernel names with
wildcard characters may be used to reduce the set of kernels. The following examples show building exactly one
or a subset of kernels for NVIDIA Ampere and Turing architecture:
### Building a subset Tensor Core GEMM kernels
To compile a subset of Tensor Core GEMM kernels with FP32 accumulation and FP16 input targetting NVIDIA Ampere and Turing architecture,
use the below cmake command line:
```bash
$ cmake .. -DCUTLASS_NVCC_ARCHS='75;80' -DCUTLASS_LIBRARY_KERNELS=cutlass_tensorop_s*gemm_f16_*_nt_align8
...
$ make cutlass_profiler -j16
```
Example command line for profiling a subset of Tensor Core GEMM kernels is as follows:
```bash
./tools/profiler/cutlass_profiler --kernels=cutlass_tensorop_s*gemm_f16_*_nt_align8 --m=3456 --n=4096 --k=4096
...
=============================
Problem ID: 1
Provider: CUTLASS
OperationKind: gemm
Operation: cutlass_tensorop_s1688gemm_f16_256x128_32x2_nt_align8
Status: Success
Verification: ON
Disposition: Passed
reference_device: Passed
cuBLAS: Passed
Arguments: --gemm_kind=universal --m=3456 --n=4096 --k=4096 --A=f16:column --B=f16:row --C=f32:column --alpha=1 \
--beta=0 --split_k_slices=1 --batch_count=1 --op_class=tensorop --accum=f32 --cta_m=256 --cta_n=128 \
--cta_k=32 --stages=2 --warps_m=4 --warps_n=2 --warps_k=1 --inst_m=16 --inst_n=8 --inst_k=8 --min_cc=75 \
--max_cc=1024
Bytes: 118489088 bytes
FLOPs: 115992428544 flops
Runtime: 1.55948 ms
Memory: 70.7616 GiB/s
Math: 74378.8 GFLOP/s
# Makefile usage
There are different sample targets for different GEMM data types and
transposititions. Be sure to specify your target architecture.
=============================
...
```
make <sgemm|dgemm|hgemm|igemm|wgemm> sm=<60|61|70> \
[transpose=<nn|nt|tn|tt>] [verbose=<0|1>] [keep=<0|1>]
### Building one CUDA Core GEMM kernel
To compile one SGEMM kernel targetting NVIDIA Ampere and Turing architecture, use the below cmake command line:
```bash
$ cmake .. -DCUTLASS_NVCC_ARCHS='75;80' -DCUTLASS_LIBRARY_KERNELS=cutlass_simt_sgemm_128x128_8x2_nn_align1
...
$ make cutlass_profiler -j16
```
Example command line for profiling single SGEMM CUDA kernel is as follows:
```bash
$ ./tools/profiler/cutlass_profiler --kernels=sgemm --m=3456 --n=4096 --k=4096
=============================
Problem ID: 1
Provider: CUTLASS
OperationKind: gemm
Operation: cutlass_simt_sgemm_128x128_8x2_nn_align1
Status: Success
Verification: ON
Disposition: Passed
cuBLAS: Passed
Arguments: --m=3456 --n=4096 --k=4096 --A=f32:column --B=f32:column --C=f32:column --alpha=1 --beta=0 --split_k_slices=1 \
--batch_count=1 --op_class=simt --accum=f32 --cta_m=128 --cta_n=128 --cta_k=8 --stages=2 --warps_m=4 \
--warps_n=2 --warps_k=1 --inst_m=1 --inst_n=1 --inst_k=1 --min_cc=50 --max_cc=1024
Bytes: 180355072 bytes
FLOPs: 115992428544 flops
Runtime: 6.73655 ms
Memory: 24.934 GiB/s
Math: 17218.4 GFLOP/s
=============================
```
### Building a subset of Tensor Core Convolution kernels
To compile a subset of Tensor core convolution kernels implementing forward propagation (fprop) with FP32 accumulation
and FP16 input targetting NVIDIA Ampere and Turing architecture, use the below cmake command line:
```bash
$ cmake .. -DCUTLASS_NVCC_ARCHS='75;80' -DCUTLASS_LIBRARY_KERNELS=cutlass_tensorop_s*fprop_optimized_f16
...
$ make cutlass_profiler -j16
```
Example command line for profiling a subset of Tensor Core convolution kernels is as follows:
```bash
$ ./tools/profiler/cutlass_profiler --kernels=cutlass_tensorop_s*fprop_optimized_f16 --n=8 --h=224 --w=224 --c=128 --k=128 --r=3 --s=3
...
=============================
Problem ID: 1
Provider: CUTLASS
OperationKind: conv2d
Operation: cutlass_tensorop_s16816fprop_optimized_f16_128x128_32x5_nhwc
Status: Success
Verification: ON
Disposition: Passed
reference_device: Passed
Arguments: --conv_kind=fprop --n=8 --h=224 --w=224 --c=128 --k=128 --r=3 --s=3 --p=224 --q=224 --pad_h=1 --pad_w=1 \
--stride_h=1 --stride_w=1 --dilation_h=1 --dilation_w=1 --Activation=f16:nhwc --Filter=f16:nhwc --Output=f32:nhwc \
--conv_mode=cross --iterator_algorithm=optimized --alpha=1 --beta=0 --split_k_mode=serial --split_k_slices=1 \
--eq_gemm_provider=none --op_class=tensorop --accum=f32 --cta_m=128 --cta_n=128 --cta_k=32 --stages=5 \
--warps_m=2 --warps_n=2 --warps_k=1 --inst_m=16 --inst_n=8 --inst_k=16 --min_cc=80 --max_cc=1024
Bytes: 1130659840 bytes
FLOPs: 118482796544 flops
Runtime: 0.711496 ms
Memory: 1479.99 GiB/s
Math: 166526 GFLOP/s
=============================
...
```
# Program usage
### Building one Convolution CUDA kernel
Program usage:
To compile and run one CUDA Core convolution kernel implementing forward propagation (fprop) with F32 accumulation
and FP32 input targetting NVIDIA Ampere and Turing architecture, use the below cmake command line:
```bash
$ cmake .. -DCUTLASS_NVCC_ARCHS='75;80' -DCUTLASS_LIBRARY_KERNELS=cutlass_simt_sfprop_optimized_128x128_8x2_nhwc
...
$ make cutlass_profiler -j16
```
<s|d|h|i|w>gemm_<nn|nt|tn|tt>
[--help]
[--schmoo || --m=<height> --n=<width> --k=<depth>]
[--i=<timing iterations>]
[--device=<device-id>]
[--alpha=<alpha> --beta=<beta>]
Example command line for profiling one CUDA Core convolution kernel:
```bash
$ ./tools/profiler/cutlass_profiler --kernels=cutlass_simt_sfprop_optimized_128x128_8x2_nhwc --n=8 --h=224 --w=224 --c=128 --k=128 --r=3 --s=3
=============================
Problem ID: 1
Provider: CUTLASS
OperationKind: conv2d
Operation: cutlass_simt_sfprop_optimized_128x128_8x2_nhwc
Status: Success
Verification: ON
Disposition: Passed
reference_device: Passed
Arguments: --conv_kind=fprop --n=8 --h=224 --w=224 --c=128 --k=128 --r=3 --s=3 --p=224 --q=224 --pad_h=1 --pad_w=1 \
--stride_h=1 --stride_w=1 --dilation_h=1 --dilation_w=1 --Activation=f32:nhwc --Filter=f32:nhwc --Output=f32:nhwc \
--conv_mode=cross --iterator_algorithm=optimized --alpha=1 --beta=0 --split_k_mode=serial --split_k_slices=1 \
--eq_gemm_provider=none --op_class=simt --accum=f32 --cta_m=128 --cta_n=128 --cta_k=8 --stages=2 --warps_m=4 \
--warps_n=2 --warps_k=1 --inst_m=1 --inst_n=1 --inst_k=1 --min_cc=50 --max_cc=1024
Bytes: 2055798784 bytes
FLOPs: 118482796544 flops
Runtime: 7.34266 ms
Memory: 260.752 GiB/s
Math: 16136.2 GFLOP/s
=============================
```
## More Details on Compiling CUTLASS Kernels and CUTLASS Profiler
- Please follow the links for more CMake examples on selectively compiling CUTLASS kernels:
- [GEMM CMake Examples](media/docs/quickstart.md#gemm-cmake-examples)
- [Implicit GEMM conovlution CMake Examples](media/docs/quickstart.md#convolution-cmake-examples)
- [Further details about the CUTLASS Profiler are described here.](media/docs/profiler.md)
# About
CUTLASS is released by NVIDIA Corporation as Open Source software under the
BSD license.
CUTLASS is released by NVIDIA Corporation as Open Source software under the
[3-clause "New" BSD license](LICENSE.txt).
# Contributors
The official list of CUTLASS developers and contributors is available here: [CONTRIBUTORS](CONTRIBUTORS.md).
# Copyright
Copyright (c) 2011-2017, NVIDIA CORPORATION. All rights reserved.
Copyright (c) 2017 - 2023 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
SPDX-License-Identifier: BSD-3-Clause
Redistribution and use in source and binary forms, with or without
modification, are not permitted.
```
Redistribution and use in source and binary forms, with or without
modification, are permitted provided that the following conditions are met:
1. Redistributions of source code must retain the above copyright notice, this
list of conditions and the following disclaimer.
2. Redistributions in binary form must reproduce the above copyright notice,
this list of conditions and the following disclaimer in the documentation
and/or other materials provided with the distribution.
3. Neither the name of the copyright holder nor the names of its
contributors may be used to endorse or promote products derived from
this software without specific prior written permission.
THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
```
THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE FOR ANY
DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
(INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
(INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.

26
bin2hex.cmake Normal file
View File

@ -0,0 +1,26 @@
# A small utility function which generates a C-header from an input file
function(FILE_TO_C_STRING FILENAME VARIABLE_NAME OUTPUT_STRING ZERO_TERMINATED)
FILE(READ "${FILENAME}" HEX_INPUT HEX)
if (${ZERO_TERMINATED})
string(APPEND HEX_INPUT "00")
endif()
string(REGEX REPLACE "(....)" "\\1\n" HEX_OUTPUT ${HEX_INPUT})
string(REGEX REPLACE "([0-9a-f][0-9a-f])" "0x\\1," HEX_OUTPUT ${HEX_OUTPUT})
set(HEX_OUTPUT "static char const ${VARIABLE_NAME}[] = {\n ${HEX_OUTPUT}\n};\n")
set(${OUTPUT_STRING} "${HEX_OUTPUT}" PARENT_SCOPE)
endfunction()
# message("Create header file for ${FILE_IN}")
# message("Create header file for ${FILE_OUT}")
file_to_c_string(${FILE_IN} ${VARIABLE_NAME} OUTPUT_STRING ZERO_TERMINATED)
set(RESULT "#pragma once\n")
string(APPEND RESULT "namespace cutlass {\n")
string(APPEND RESULT "namespace nvrtc {\n")
string(APPEND RESULT "${OUTPUT_STRING}")
string(APPEND RESULT "} // namespace nvrtc\n")
string(APPEND RESULT "} // namespace cutlass\n")
file(WRITE "${FILE_OUT}" "${RESULT}")

View File

@ -0,0 +1,21 @@
# Generated file
if (DEFINED ENV{CUTLASS_TEST_EXECUTION_ENVIRONMENT})
set(_CUTLASS_TEST_EXECUTION_ENVIRONMENT $ENV{CUTLASS_TEST_EXECUTION_ENVIRONMENT})
else()
set(_CUTLASS_TEST_EXECUTION_ENVIRONMENT @CUTLASS_TEST_EXECUTION_ENVIRONMENT@)
endif()
if (NOT "@TEST_EXE_DIR@" STREQUAL "")
set(TEST_EXE_PATH @TEST_EXE_DIR@/@TEST_EXE@)
else()
set(TEST_EXE_PATH @TEST_EXE@)
endif()
add_test("@TEST_NAME@" ${_CUTLASS_TEST_EXECUTION_ENVIRONMENT} "${TEST_EXE_PATH}" @TEST_COMMAND_OPTIONS@)
if (NOT "@TEST_EXE_WORKING_DIRECTORY@" STREQUAL "")
set_tests_properties("@TEST_NAME@" PROPERTIES WORKING_DIRECTORY "@TEST_EXE_WORKING_DIRECTORY@")
endif()
set_tests_properties(@TEST_NAME@ PROPERTIES DISABLED @__DISABLE_TESTS@)

View File

@ -0,0 +1,7 @@
get_filename_component(NvidiaCutlass_CMAKE_DIR "${CMAKE_CURRENT_LIST_FILE}" PATH)
include(CMakeFindDependencyMacro)
if(NOT TARGET nvidia::cutlass::CUTLASS)
include("${NvidiaCutlass_CMAKE_DIR}/NvidiaCutlassTargets.cmake")
endif()

View File

@ -0,0 +1,14 @@
set(CPACK_PACKAGE_NAME NvidiaCutlass)
set(CPACK_PACKAGE_VENDOR NVIDIA)
set(CPACK_PACKAGE_CONTACT info@nvidia.com)
set(CPACK_PACKAGE_DESCRIPTION_SUMMARY "CUTLASS CUDA C++ Template Linear Algebra Library")
set(CPACK_PACKAGE_INSTALL_DIRECTORY ${CPACK_PACKAGE_NAME})
set(CPACK_PACKAGE_VERSION_MAJOR ${PROJECT_VERSION_MAJOR})
set(CPACK_PACKAGE_VERSION_MINOR ${PROJECT_VERSION_MINOR})
set(CPACK_PACKAGE_VERSION_PATCH ${PROJECT_VERSION_PATCH})
set(CPACK_VERBATIM_VARIABLES YES)
# set(CPACK_PACKAGE_DESCRIPTION_FILE ${CMAKE_CURRENT_LIST_DIR}/Description.txt)
# set(CPACK_RESOURCE_FILE_WELCOME ${CMAKE_CURRENT_LIST_DIR}/Welcome.txt)
# set(CPACK_RESOURCE_FILE_LICENSE ${CMAKE_CURRENT_LIST_DIR}/License.txt)
# set(CPACK_RESOURCE_FILE_README ${CMAKE_CURRENT_LIST_DIR}/Readme.txt)
include(CPack)

23
cmake/googletest.cmake Normal file
View File

@ -0,0 +1,23 @@
include(FetchContent)
set(GOOGLETEST_DIR "" CACHE STRING "Location of local GoogleTest repo to build against")
if(GOOGLETEST_DIR)
set(FETCHCONTENT_SOURCE_DIR_GOOGLETEST ${GOOGLETEST_DIR} CACHE STRING "GoogleTest source directory override")
endif()
FetchContent_Declare(
googletest
GIT_REPOSITORY https://github.com/google/googletest.git
GIT_TAG 0fe9660
)
FetchContent_GetProperties(googletest)
if(NOT googletest_POPULATED)
FetchContent_Populate(googletest)
if (MSVC)
set(gtest_force_shared_crt ON CACHE BOOL "" FORCE)
endif()
add_subdirectory(${googletest_SOURCE_DIR} ${googletest_BINARY_DIR} EXCLUDE_FROM_ALL)
endif()

49
cmake/nop.cu Normal file
View File

@ -0,0 +1,49 @@
/***************************************************************************************************
* Copyright (c) 2017 - 2023 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
* SPDX-License-Identifier: BSD-3-Clause
*
* Redistribution and use in source and binary forms, with or without
* modification, are permitted provided that the following conditions are met:
*
* 1. Redistributions of source code must retain the above copyright notice, this
* list of conditions and the following disclaimer.
*
* 2. Redistributions in binary form must reproduce the above copyright notice,
* this list of conditions and the following disclaimer in the documentation
* and/or other materials provided with the distribution.
*
* 3. Neither the name of the copyright holder nor the names of its
* contributors may be used to endorse or promote products derived from
* this software without specific prior written permission.
*
* THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
* AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
* IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
* DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
* FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
* DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
* SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
* CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
* OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
* OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
*
**************************************************************************************************/
/*! \file
\brief Basic CUDA file for testing compiler flags.
*/
__device__ int inner()
{
return -1;
}
__global__ void test()
{
inner();
}
int main()
{
test<<<1,1>>>();
return 0;
}

38
cmake/version.h.in Normal file
View File

@ -0,0 +1,38 @@
#include <cstdint>
#include <string>
#define CUTLASS_MAJOR @CUTLASS_VERSION_MAJOR@
#define CUTLASS_MINOR @CUTLASS_VERSION_MINOR@
#define CUTLASS_PATCH @CUTLASS_VERSION_PATCH@
#define CUTLASS_BUILD @CUTLASS_VERSION_BUILD@
#define CUTLASS_VERSION ((CUTLASS_MAJOR)*100 + (CUTLASS_MINOR)*10 + CUTLASS_PATCH)
namespace cutlass {
inline uint32_t getVersion() {
return CUTLASS_VERSION;
}
inline uint32_t getVersionMajor() {
return CUTLASS_MAJOR;
}
inline uint32_t getVersionMinor() {
return CUTLASS_MINOR;
}
inline uint32_t getVersionPatch() {
return CUTLASS_PATCH;
}
inline uint32_t getVersionBuild() {
return CUTLASS_BUILD + 0;
}
inline std::string getVersionString() {
std::string version = "@CUTLASS_VERSION@";
if (getVersionBuild()) {
version += "." + std::to_string(getVersionBuild());
}
return version;
}
inline std::string getGitRevision() {
return "@CUTLASS_REVISION@";
}
} // namespace cutlass

154
common.mk
View File

@ -1,154 +0,0 @@
#/******************************************************************************
# * Copyright (c) 2011, Duane Merrill. All rights reserved.
# * Copyright (c) 2011-2017, NVIDIA CORPORATION. All rights reserved.
# *
# * Redistribution and use in source and binary forms, with or without
# * modification, are permitted provided that the following conditions are met:
# * * Redistributions of source code must retain the above copyright
# * notice, this list of conditions and the following disclaimer.
# * * Redistributions in binary form must reproduce the above copyright
# * notice, this list of conditions and the following disclaimer in the
# * documentation and/or other materials provided with the distribution.
# * * Neither the name of the NVIDIA CORPORATION nor the
# * names of its contributors may be used to endorse or promote products
# * derived from this software without specific prior written permission.
# *
# * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
# * ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
# * WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
# * DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE FOR ANY
# * DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
# * (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
# * LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
# * ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
# * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
# * SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
# *
#******************************************************************************/
#-------------------------------------------------------------------------------
# Commandline Options
#-------------------------------------------------------------------------------
# sm=<XX,...> Compute-capability to compile for, e.g., "sm=200,300,350" (SM2.0 by default).
COMMA := ,
ifdef sm
SM_ARCH := $(subst $(COMMA),-,$(sm))
else
$(error Please specify SM architecture makefile argument: "sm=XX")
endif
ifeq (70, $(findstring 70, $(SM_ARCH)))
SM_TARGETS += -gencode=arch=compute_70,code=\"sm_70,compute_70\"
endif
ifeq (62, $(findstring 62, $(SM_ARCH)))
SM_TARGETS += -gencode=arch=compute_62,code=\"sm_62,compute_62\"
endif
ifeq (61, $(findstring 61, $(SM_ARCH)))
SM_TARGETS += -gencode=arch=compute_61,code=\"sm_61,compute_61\"
endif
ifeq (60, $(findstring 60, $(SM_ARCH)))
SM_TARGETS += -gencode=arch=compute_60,code=\"sm_60,compute_60\"
endif
ifeq (52, $(findstring 52, $(SM_ARCH)))
SM_TARGETS += -gencode=arch=compute_52,code=\"sm_52,compute_52\"
endif
ifeq (37, $(findstring 37, $(SM_ARCH)))
SM_TARGETS += -gencode=arch=compute_37,code=\"sm_37,compute_37\"
endif
ifeq (35, $(findstring 35, $(SM_ARCH)))
SM_TARGETS += -gencode=arch=compute_35,code=\"sm_35,compute_35\"
endif
ifeq (30, $(findstring 30, $(SM_ARCH)))
SM_TARGETS += -gencode=arch=compute_30,code=\"sm_30,compute_30\"
endif
ifeq (21, $(findstring 21, $(SM_ARCH)))
SM_TARGETS += -gencode=arch=compute_20,code=\"sm_21,compute_20\"
endif
ifeq (20, $(findstring 20, $(SM_ARCH)))
SM_TARGETS += -gencode=arch=compute_20,code=\"sm_20,compute_20\"
endif
# [verbose=<0|1>] Verbose toolchain output from nvcc option
ifeq ($(verbose), 1)
NVCCFLAGS += -v
endif
# [keep=<0|1>] Keep intermediate compilation artifacts option
ifeq ($(keep), 1)
NVCCFLAGS += -keep
endif
# [debug=<0|1>] Generate debug mode code
ifeq ($(debug), 1)
NVCCFLAGS += -G
endif
#-------------------------------------------------------------------------------
# Compiler and compilation platform
#-------------------------------------------------------------------------------
BASE_DIR := $(dir $(lastword $(MAKEFILE_LIST)))
NVCC := "$(shell which nvcc)"
ifdef nvccver
NVCC_VERSION := $(nvccver)
else
NVCC_VERSION := $(strip $(shell nvcc --version | grep release | sed 's/.*release //' | sed 's/,.*//'))
endif
# Detect OS
OSUPPER := $(shell uname -s 2>/dev/null | tr [:lower:] [:upper:])
# Default flags: verbose kernel properties (regs, smem, cmem, etc.); runtimes for compilation phases
NVCCFLAGS += -O3 -Xptxas -v
ifeq (WIN_NT, $(findstring WIN_NT, $(OSUPPER)))
# For MSVC
# Enable more warnings and treat as errors
NVCCFLAGS += -Xcompiler /W3 -Xcompiler /WX
# Disable excess x86 floating point precision that can lead to results being labeled incorrectly
NVCCFLAGS += -Xcompiler /fp:strict
# Compiler
CC := cl
# Multithreaded runtime
NVCCFLAGS += -Xcompiler /MT
CUDART_CYG := "$(shell dirname $(NVCC))/../lib/x64/cudart.lib"
CUDART := "$(shell cygpath -w $(CUDART_CYG))"
else
# For g++
# Disable excess x86 floating point precision that can lead to results being labeled incorrectly
#NVCCFLAGS += -Xcompiler -ffloat-store
# Compiler
CC := g++
CUDART := "$(shell dirname $(NVCC))/../lib64/libcudart_static.a"
endif
# Suffix to append to each binary
BIN_SUFFIX := sm$(SM_ARCH)_$(NVCC_VERSION)
#-------------------------------------------------------------------------------
# Function for computing dependency Lists
#-------------------------------------------------------------------------------
rwildcard=$(foreach d,$(wildcard $1*),$(call rwildcard,$d/,$2) $(filter $(subst *,%,$2),$d))

152
cuBLAS.cmake Normal file
View File

@ -0,0 +1,152 @@
# Copyright (c) 2017 - 2023 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: BSD-3-Clause
#
# Redistribution and use in source and binary forms, with or without
# modification, are permitted provided that the following conditions are met:
#
# 1. Redistributions of source code must retain the above copyright notice, this
# list of conditions and the following disclaimer.
#
# 2. Redistributions in binary form must reproduce the above copyright notice,
# this list of conditions and the following disclaimer in the documentation
# and/or other materials provided with the distribution.
#
# 3. Neither the name of the copyright holder nor the names of its
# contributors may be used to endorse or promote products derived from
# this software without specific prior written permission.
#
# THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
# AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
# IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
# DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
# FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
# DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
# SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
# CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
# OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
# OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
message(STATUS "Configuring cublas ...")
if((DEFINED CUTLASS_ENABLE_CUBLAS AND NOT CUTLASS_ENABLE_CUBLAS) OR
(DEFINED CUBLAS_ENABLED AND NOT CUBLAS_ENABLED))
# Don't add cuBLAS if it's defined and false, assume it's not found.
set(CUBLAS_FOUND OFF)
message(STATUS "cuBLAS Disabled.")
elseif(NOT TARGET cublas)
find_path(
_CUBLAS_INCLUDE_DIR
NAMES cublas_v2.h
HINTS
${CUBLAS_INCLUDE_PATH}
ENV CUBLAS_INCLUDE_PATH
${CUBLAS_PATH}
ENV CUBLAS_PATH
${CUDA_TOOLKIT_ROOT_DIR}
PATH_SUFFIXES
include
)
find_library(
_CUBLAS_LIBRARY
NAMES cublas
HINTS
${CUBLAS_LIBRARY_PATH}
ENV CUBLAS_LIBRARY_PATH
${_CUBLAS_INCLUDE_DIR}/..
${CUBLAS_PATH}
ENV CUBLAS_PATH
${CUDA_TOOLKIT_ROOT_DIR}
PATH_SUFFIXES
lib64
lib/x64
lib
)
if(_CUBLAS_INCLUDE_DIR AND _CUBLAS_LIBRARY)
message(STATUS "cuBLAS: ${_CUBLAS_LIBRARY}")
message(STATUS "cuBLAS: ${_CUBLAS_INCLUDE_DIR}")
set(CUBLAS_FOUND ON CACHE INTERNAL "cublas Library Found")
set(CUBLAS_LIBRARY ${_CUBLAS_LIBRARY})
set(CUBLAS_INCLUDE_DIR ${_CUBLAS_INCLUDE_DIR})
else()
message(STATUS "cublas not found.")
set(CUBLAS_FOUND OFF CACHE INTERNAL "cublas Library Found")
endif()
endif()
set(CUTLASS_ENABLE_CUBLAS ${CUBLAS_FOUND} CACHE BOOL "Enable CUTLASS to build with cuBLAS library.")
if(CUTLASS_ENABLE_CUBLAS AND NOT CUBLAS_FOUND)
message(FATAL_ERROR "CUTLASS_ENABLE_CUBLAS enabled but cuBLAS library could not be found.")
endif()
if(CUTLASS_ENABLE_CUBLAS AND NOT TARGET cublas)
if(WIN32)
add_library(cublas STATIC IMPORTED GLOBAL)
else()
add_library(cublas SHARED IMPORTED GLOBAL)
endif()
add_library(nvidia::cublas ALIAS cublas)
set_property(
TARGET cublas
PROPERTY IMPORTED_LOCATION
${CUBLAS_LIBRARY})
target_include_directories(
cublas
INTERFACE
$<INSTALL_INTERFACE:include>
$<BUILD_INTERFACE:${CUBLAS_INCLUDE_DIR}>)
find_library(
_CUBLASLT_LIBRARY
NAMES cublasLt
HINTS
${CUBLAS_LIBRARY_PATH}
ENV CUBLAS_LIBRARY_PATH
${_CUBLAS_INCLUDE_DIR}/..
${CUBLAS_PATH}
ENV CUBLAS_PATH
${CUDA_TOOLKIT_ROOT_DIR}
PATH_SUFFIXES
lib64
lib/x64
lib
)
if(_CUBLASLT_LIBRARY AND NOT TARGET cublasLt)
if(WIN32)
add_library(cublasLt STATIC IMPORTED GLOBAL)
else()
add_library(cublasLt SHARED IMPORTED GLOBAL)
endif()
set_property(
TARGET cublasLt
PROPERTY IMPORTED_LOCATION
${_CUBLASLT_LIBRARY})
add_library(nvidia::cublasLt ALIAS cublasLt)
target_link_libraries(cublas INTERFACE cublasLt)
endif()
endif()
message(STATUS "Configuring cuBLAS ... done.")

112
cuDNN.cmake Normal file
View File

@ -0,0 +1,112 @@
# Copyright (c) 2017 - 2023 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
# SPDX-License-Identifier: BSD-3-Clause
#
# Redistribution and use in source and binary forms, with or without
# modification, are permitted provided that the following conditions are met:
#
# 1. Redistributions of source code must retain the above copyright notice, this
# list of conditions and the following disclaimer.
#
# 2. Redistributions in binary form must reproduce the above copyright notice,
# this list of conditions and the following disclaimer in the documentation
# and/or other materials provided with the distribution.
#
# 3. Neither the name of the copyright holder nor the names of its
# contributors may be used to endorse or promote products derived from
# this software without specific prior written permission.
#
# THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
# AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
# IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
# DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
# FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
# DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
# SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
# CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
# OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
# OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
if(DEFINED CUDNN_ENABLED)
set(CUTLASS_ENABLE_CUDNN ${CUDNN_ENABLED} CACHE BOOL "Enable CUTLASS to build with cuDNN library.")
endif()
if(DEFINED CUTLASS_ENABLE_CUDNN AND NOT CUTLASS_ENABLE_CUDNN)
return()
endif()
message(STATUS "Configuring cuDNN ...")
find_path(
_CUDNN_INCLUDE_DIR cudnn.h
PATHS
${CUDA_TOOLKIT_ROOT_DIR}/include
$ENV{CUDNN_PATH}/include
$ENV{CUDA_PATH}/include
${CUDNN_PATH}/include
/usr/include)
find_library(
_CUDNN_LIBRARY cudnn
HINTS
${CUDA_TOOLKIT_ROOT_DIR}/lib64
${CUDA_TOOLKIT_ROOT_DIR}/lib/x64
${CUDA_TOOLKIT_ROOT_DIR}/lib
$ENV{CUDNN_PATH}/lib64
$ENV{CUDNN_PATH}/lib/x64
$ENV{CUDNN_PATH}/lib
$ENV{CUDA_PATH}/lib64
$ENV{CUDA_PATH}/lib/x64
$ENV{CUDA_PATH}/lib
${CUDNN_PATH}/lib64
${CUDNN_PATH}/lib/x64
${CUDNN_PATH}/lib
/usr/lib/x86_64-linux-gnu
/usr/lib)
if(_CUDNN_INCLUDE_DIR AND _CUDNN_LIBRARY)
message(STATUS "cuDNN: ${_CUDNN_LIBRARY}")
message(STATUS "cuDNN: ${_CUDNN_INCLUDE_DIR}")
set(CUDNN_FOUND ON CACHE INTERNAL "cuDNN Library Found")
else()
message(STATUS "cuDNN not found.")
set(CUDNN_FOUND OFF CACHE INTERNAL "cuDNN Library Found")
endif()
set(CUTLASS_ENABLE_CUDNN ${CUDNN_FOUND} CACHE BOOL "Enable CUTLASS to build with cuDNN library.")
if (CUTLASS_ENABLE_CUDNN AND NOT TARGET cudnn)
set(CUDNN_INCLUDE_DIR ${_CUDNN_INCLUDE_DIR})
set(CUDNN_LIBRARY ${_CUDNN_LIBRARY})
if(WIN32)
add_library(cudnn STATIC IMPORTED GLOBAL)
else()
add_library(cudnn SHARED IMPORTED GLOBAL)
endif()
add_library(nvidia::cudnn ALIAS cudnn)
set_property(
TARGET cudnn
PROPERTY IMPORTED_LOCATION
${CUDNN_LIBRARY})
target_include_directories(
cudnn
INTERFACE
$<INSTALL_INTERFACE:include>
$<BUILD_INTERFACE:${CUDNN_INCLUDE_DIR}>)
endif()
if(CUTLASS_ENABLE_CUDNN AND NOT CUDNN_FOUND)
message(FATAL_ERROR "CUTLASS_ENABLE_CUDNN enabled but cuDNN library could not be found.")
endif()
message(STATUS "Configuring cuDNN ... done.")

View File

@ -1,154 +0,0 @@
/******************************************************************************
* Copyright (c) 2011-2017, NVIDIA CORPORATION. All rights reserved.
*
* Redistribution and use in source and binary forms, with or without
* modification, are not permitted.
*
* THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
* ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
* WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
* DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE FOR ANY
* DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
* (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
* LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
* ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
* (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
* SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
*
******************************************************************************/
#pragma once
/**
* \file
* block-wide tile-loading abstractions
*/
#include "../util/util.h"
namespace cutlass {
namespace gemm {
/******************************************************************************
* load_algorithm
******************************************************************************/
/**
* \brief Enumeration of matrix loading algorithms
*/
struct load_algorithm
{
/// \brief Enumerants. See corresponding tag types.
enum kind_t
{
CongruousCopy = 0,
CrosswiseCopy = 1,
};
/**
* \brief Generic tag
*/
template <kind_t Kind>
struct any_tag : nv_std::integral_constant<kind_t, Kind> {};
/**
* \brief Copy from a global matrix that is row-major in relation
* to the local row-major tile
*/
typedef any_tag<CongruousCopy> contiguous_tag_t;
/**
* \brief Copy from a global matrix that is column-major in relation
* to the local row-major tile
*/
typedef any_tag<CrosswiseCopy> crosswise_tag_t;
};
/******************************************************************************
* block_loader
******************************************************************************/
/**
* \brief A three-phase data loading abstraction (prefetch, commit, and
* advance) for iterating over ranges of block-wide matrix tiles.
*
* Each iteration sequence produces a KxL (height-by-width) block-wide tile of
* value_t in shared memory. The layout of the shared
* block-wide tile is a row-major (L-major) tiling of dp_vector_t items, which are
* themselves column-major (K-major) vectors of value_t. Its dimensions are:
* K = BlockDpVectorsK * (sizeof(dp_vector_t) / sizeof(value_t)
* L = BlockDpVectorsL
*
* NB: This generic class is not directly constructible. Architecture- and
* algorithm-specific template specializations will provide the API
* functionality prescribed here.
*
*/
template <
int BlockThreads, ///< Number of threads in each thread block (blockDim.x)
int BlockDpVectorsK, ///< Extent of block-wide tile in dp_vector_t along the K-axis (height)
int BlockDpVectorsL, ///< Extent of block-wide tile in dp_vector_t along the L-axis (width)
typename value_t, ///< Input matrix value type
int LeadingDimAlignBytes, ///< Byte alignment of input matrix leading dimension
bool AllowRaggedTiles, ///< Whether the input matrix's dimensions need not be an even-multiple of the block-wide tile dimensions
typename dp_vector_t, ///< Dot-product vector type along the K-axis
load_algorithm::kind_t LoadAlgorithm> ///< Algorithm for loading a shared tile of KxL matrix data
struct block_loader
{
//-------------------------------------------------------------------------
// Constructor API
//-------------------------------------------------------------------------
/// Constructor
block_loader(
value_t *d_matrix, ///< Pointer to input matrix
int matrix_values_l, ///< Extent of the input matrix in value_t along the L-axis
int matrix_values_stride_k, ///< Distance in value_t within pitched-linear memory between successive coordinates along the K-axis
int matrix_values_stride_l, ///< Distance in value_t within pitched-linear memory between successive coordinates along the L-axis
int2 block_begin_item_coords, ///< Thread block's starting value_t coordinates (l, k) within the input matrix
int block_end_item_k); ///< Thread block's ending coordinate (k) within the input matrix (one-past)
//-------------------------------------------------------------------------
// Loader API
//-------------------------------------------------------------------------
/**
* Request the current block-wide tile
*/
void request();
/**
* Advance the loader to the next block-wide tile in the K-axis
*/
void next();
/**
* Commit the previously-requested block-wide tile to shared memory
*
* NB: To facilitate padding for avoiding shared memory bank conflicts, we
* allow the row stride _BlockDpVectorsL to be arbitrarily bigger than the
* tile width BlockDpVectorsL.
*/
template <int _BlockDpVectorsL>
void commit(
dp_vector_t (&scratch_tile)[BlockDpVectorsK][_BlockDpVectorsL]);
};
} // namespace gemm
} // namespace cutlass
/******************************************************************************
* Tail-include specializations that adhere to the block_loader API
******************************************************************************/
#include "block_loader_crosswise.h"
#include "block_loader_congruous_dp1.h"
#include "block_loader_congruous_idp4.h"

View File

@ -1,398 +0,0 @@
/******************************************************************************
* Copyright (c) 2011-2017, NVIDIA CORPORATION. All rights reserved.
*
* Redistribution and use in source and binary forms, with or without
* modification, are not permitted.
*
* THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
* ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
* WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
* DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE FOR ANY
* DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
* (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
* LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
* ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
* (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
* SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
*
******************************************************************************/
#pragma once
/**
* \file
* Tile-loading abstraction for thread blocks
*/
#include "../util/util.h"
namespace cutlass {
namespace gemm {
/******************************************************************************
* block_loader (CongruousCopy + dp1 specialization)
******************************************************************************/
/**
* \brief A three-phase data loading abstraction (prefetch, commit, and
* advance) for iterating over ranges of block-wide matrix tiles.
* (CongruousCopy + dp1 specialization)
*
* Each iteration sequence produces a KxL (height-by-width) block-wide tile of
* value_t in shared memory. The layout of the shared block-wide tile is
* a row-major (L-major) tiling of singleton "dp1" dp_vector_t items, where
* dp_vector_t == value_t. Its dimensions are:
* K = BlockDpVectorsK
* L = BlockDpVectorsL
*
* The data is copied from a corresponding tile of global matrix data whose
* layout of value_t is also L-major. This constitutes a CongruousCopy
* between the L-major global tile and the L-major shared tile.
*
* NB: Because they are "dp1" singletons, the K-major orientation of
* dp_vector_t in shared memory is irrelevant, and the L-major global and
* shared tile layouts are perfectly congruous. As a result, we can increase
* the granularity of data transfer via vectorization of loads and stores
* without any intermediate {dis|re}assembly.
*
* NB: Consecutive threads within a block are mapped in L-major
* fashion across a first-set of LDG-vectors of dp_vector_t (value_t) within
* their global tile. Successive sets of LDG-vectors are then strip-mined
* as necessary down the K-axis. These discontiguous LDG-vectors comprise the
* thread's "slice" of the block-wide tile.
*/
template <
int BlockThreads, ///< Number of threads in each thread block (blockDim.x)
int BlockDpVectorsK, ///< Extent of block-wide tile in dp_vector_t along the K-axis (height)
int BlockDpVectorsL, ///< Extent of block-wide tile in dp_vector_t along the L-axis (width)
typename value_t, ///< Input matrix value type
int LeadingDimAlignBytes, ///< Byte alignment of input matrix leading dimension
bool AllowRaggedTiles ///< Whether the input matrix's dimensions need not be an even-multiple of the block-wide tile dimensions
>
struct block_loader<
BlockThreads,
BlockDpVectorsK,
BlockDpVectorsL,
value_t,
LeadingDimAlignBytes,
AllowRaggedTiles,
value_t, ///< Dot-product vector type along the K-axis (dp1 specialization)
load_algorithm::CongruousCopy> ///< Algorithm for loading a shared tile of KxL matrix data (CongruousCopy specialization)
{
//-------------------------------------------------------------------------
// Constants and types
//-------------------------------------------------------------------------
/// Dot-product vector type along the K-axis
typedef value_t dp_vector_t;
enum
{
/// Number of value_t in a dp_vector_t
DpVectorItems = divide_assert<sizeof(dp_vector_t), sizeof(value_t)>::value,
/// Number of dp_vector_t in a block-wide tile
BlockDpVectors = BlockDpVectorsK * BlockDpVectorsL,
/// Number of dp_vector_t in a thread-tile
ThreadDpVectors = divide_assert<BlockDpVectors, BlockThreads>::value,
};
/// Data movement type, coarsened by LeadingDimAlignBytes, capped by the
/// smaller of either ThreadDpVectors or BlockDpVectorsL
typedef io_vector<
dp_vector_t,
__NV_STD_MIN(ThreadDpVectors, BlockDpVectorsL),
LeadingDimAlignBytes>
ldg_vector_t;
enum
{
/// Number of dp_vector_t per ldg_vector_t
LdgVectorDpVectors = ldg_vector_t::VectorItems,
/// Number of value_t per ldg_vector_t
LdgVectorItems = LdgVectorDpVectors * DpVectorItems,
/// Total number of ldg_vector_t within each block-wide tile
BlockLdgVectors = divide_assert<BlockDpVectors, LdgVectorDpVectors>::value,
/// Extent of the block-wide tile in ldg_vector_t along L-axis
BlockLdgVectorsL = divide_assert<BlockDpVectorsL, LdgVectorDpVectors>::value,
/// Extent of the block-wide tile in ldg_vector_t along K-axis
BlockLdgVectorsK = BlockDpVectorsK,
/// Number of ldg_vector_t within each thread-tile
ThreadLdgVectors = divide_assert<BlockLdgVectors, BlockThreads>::value,
/// Extent of the thread tile in ldg_vector_t along L-axis
ThreadLdgVectorsL = __NV_STD_MAX(1, (BlockLdgVectorsL / BlockThreads)),
/// Extent of the thread tile in ldg_vector_t along K-axis
ThreadLdgVectorsK = divide_assert<ThreadLdgVectors, ThreadLdgVectorsL>::value,
/// Number of ldg_vector_t within each stripmine-tile
StripmineLdgVectors = BlockThreads,
/// Extent of the stripmine tile in ldg_vector_t along L-axis
StripmineLdgVectorsL = __NV_STD_MIN(BlockLdgVectorsL, StripmineLdgVectors),
/// Extent of the stripmine tile in ldg_vector_t along K-axis
StripmineLdgVectorsK = divide_assert<StripmineLdgVectors, StripmineLdgVectorsL>::value,
/// Alignment in dp_vector_t along L needed for committing prefetch
AlignmentDpVectorsL = LdgVectorDpVectors,
};
/// Predicate bit vector
typedef uint64_t predicate_mask_t;
//-------------------------------------------------------------------------
// Assert assumptions
//-------------------------------------------------------------------------
static_assert(
(ThreadLdgVectors <= sizeof(predicate_mask_t) * 8),
"Predicate mask type does not contain enough bits for encoding load predicates");
//-------------------------------------------------------------------------
// Members
//-------------------------------------------------------------------------
/// Input pointer to matrix in ldg_vector_t
ldg_vector_t *d_matrix_ldgvecs;
/// Extent of the input matrix in ldg_vector_t along the L-axis
int matrix_ldgvecs_l;
/// Thread block's ending ldg_vector_t coordinate (k) within the input matrix (one-past)
int block_end_ldgvec_k;
/// Predicate bits for guarding ldg_vector_t loads within "whole-k" block-wide tiles
predicate_mask_t guard;
/// Predicate bits for guarding ldg_vector_t loads within the final block-wide "residue" tile
predicate_mask_t residue_guard;
/// Iteration span in "whole-k" block-wide tiles
int wholek_tiles_remaining;
/// Distance in ldg_vector_t within pitched-linear memory between successive coordinates along the K-axis
int matrix_ldgvec_stride_k;
/// Distance in ldg_vector_t within pitched-linear memory between successive coordinates along the L-axis
int matrix_ldgvec_stride_l;
/// ldg_vector_t coordinates (l, k) of thread-tile within the block-wide tile
int2 block_thread_ldgvec_coords;
/// Thread-wide tile of prefetch data
ldg_vector_t thread_tile[ThreadLdgVectorsK][ThreadLdgVectorsL];
//-------------------------------------------------------------------------
// Constructor API
//-------------------------------------------------------------------------
/// Constructor
inline __device__
block_loader(
value_t *d_matrix_items, ///< Input pointer to matrix in value_t
int matrix_items_l, ///< Extent of the input matrix in value_t along the L-axis
int matrix_items_stride_k, ///< Distance in value_t within pitched-linear memory between successive coordinates along the K-axis
int matrix_items_stride_l, ///< Distance in value_t within pitched-linear memory between successive coordinates along the L-axis
int2 matrix_block_item_coords, ///< value_t coordinates (l, k) of first block-wide tile within the input matrix
int block_end_item_k) ///< Thread block's ending coordinate (k) within the input matrix (one-past)
:
block_end_ldgvec_k(block_end_item_k),
guard(0),
residue_guard(0)
{
matrix_ldgvecs_l = matrix_items_l / LdgVectorItems;
matrix_ldgvec_stride_k = matrix_items_stride_k / LdgVectorItems,
matrix_ldgvec_stride_l = matrix_items_stride_l;
// ldg_vector_t coordinates (l, k) of thread-tile within the block-wide tile
block_thread_ldgvec_coords = make_int2(
threadIdx.x % BlockLdgVectorsL, // l-coordinate
threadIdx.x / BlockLdgVectorsL); // k-coordinate
// ldg_vector_t coordinates (l, k) of first block-wide tile within the input matrix
int2 matrix_block_ldgvec_coords = make_int2(
matrix_block_item_coords.x / LdgVectorItems, // l-coordinate
matrix_block_item_coords.y); // k-coordinate
// Iteration span in ldg_vector_t
int span_ldgvec_k = (block_end_item_k - matrix_block_item_coords.y);
// ldg_vector_t coordinates (l, k) of first thread-tile tile within the input matrix
int2 matrix_thread_ldgvec_coords = make_int2(
block_thread_ldgvec_coords.x + matrix_block_ldgvec_coords.x,
block_thread_ldgvec_coords.y + matrix_block_ldgvec_coords.y);
// Iteration range in "whole-k" block-wide tiles
wholek_tiles_remaining = span_ldgvec_k / BlockLdgVectorsK;
// Extent of final residue-tile in ldg_vector_t along K-axis
int residue_ldgvecs_k = span_ldgvec_k % BlockLdgVectorsK;
// Initialize I/O predicates
if (AllowRaggedTiles)
{
// Outer thread-tile ldg_vector_t iteration (K-axis)
#pragma unroll
for (int thread_ldgvec_k = 0; thread_ldgvec_k < ThreadLdgVectorsK; ++thread_ldgvec_k)
{
int block_ldgvec_k = block_thread_ldgvec_coords.y + (thread_ldgvec_k * StripmineLdgVectorsK);
// Whether block_ldgvec_coords.y is valid in the final residue tile
predicate_mask_t valid_k = (block_ldgvec_k < residue_ldgvecs_k);
// Inner thread-tile ldg_vector_t iteration (L-axis)
#pragma unroll
for (int thread_ldgvec_l = 0; thread_ldgvec_l < ThreadLdgVectorsL; ++thread_ldgvec_l)
{
int block_ldgvec_l = block_thread_ldgvec_coords.x + (thread_ldgvec_l * StripmineLdgVectorsL);
// Whether block_ldgvec_coords.x is valid any block-wide tile
predicate_mask_t valid_l = (matrix_block_ldgvec_coords.x + block_ldgvec_l < matrix_ldgvecs_l);
// Linear index of ldg_vector_t load
int ldgvec_idx = thread_ldgvec_l + (thread_ldgvec_k * ThreadLdgVectorsL);
// Set predicate guard bits
guard |= (valid_l << ldgvec_idx);
residue_guard |= ((valid_l & valid_k) << ldgvec_idx);
}
}
// Promote residue-guard to primary-guard if no full tiles remain
if (!wholek_tiles_remaining)
{
guard = residue_guard;
}
}
// Update the input pointer to be matrix_thread_ldgvec_coords
this->d_matrix_ldgvecs =
reinterpret_cast<ldg_vector_t*>(d_matrix_items) +
(matrix_thread_ldgvec_coords.y * matrix_ldgvec_stride_k) +
(matrix_thread_ldgvec_coords.x * matrix_ldgvec_stride_l);
}
//-------------------------------------------------------------------------
// Loader API
//-------------------------------------------------------------------------
/**
* Request the current block-wide tile
*/
inline __device__
void request()
{
// Outer thread-tile ldg_vector_t iteration (K-axis)
#pragma unroll
for (int thread_ldgvec_k = 0; thread_ldgvec_k < ThreadLdgVectorsK; ++thread_ldgvec_k)
{
// Inner thread-tile ldg_vector_t iteration (L-axis)
#pragma unroll
for (int thread_ldgvec_l = 0; thread_ldgvec_l < ThreadLdgVectorsL; ++thread_ldgvec_l)
{
// Linear index of ldg_vector_t load
int ldgvec_idx = (thread_ldgvec_k * ThreadLdgVectorsL) + thread_ldgvec_l;
// Unpack predicate guard
predicate_mask_t valid = ((guard >> ldgvec_idx) & 1);
if (!AllowRaggedTiles || valid)
{
// Perform load
thread_tile[thread_ldgvec_k][thread_ldgvec_l].load(
d_matrix_ldgvecs +
(thread_ldgvec_k * StripmineLdgVectorsK * matrix_ldgvec_stride_k) +
(thread_ldgvec_l * StripmineLdgVectorsL * matrix_ldgvec_stride_l));
}
else
{
// Zero-initialize
#pragma unroll
for (int dpvec = 0; dpvec < LdgVectorDpVectors; ++dpvec)
thread_tile[thread_ldgvec_k][thread_ldgvec_l].buff[dpvec] = 0;
}
}
}
}
/**
* Advance the loader to the next block-wide tile in the K-axis
*/
inline __device__
void next()
{
d_matrix_ldgvecs += (matrix_ldgvec_stride_k * BlockLdgVectorsK);
if (AllowRaggedTiles)
{
--wholek_tiles_remaining;
// Promote residue-guard to primary-guard if no full tiles remain
if (!wholek_tiles_remaining)
{
guard = residue_guard;
}
}
}
/**
* Commit the previously-requested block-wide tile to shared memory
*
* NB: To facilitate padding for avoiding shared memory bank conflicts, we
* allow the row stride SmemDpVectorsL to be arbitrarily bigger than the
* tile width BlockDpVectorsL.
*/
template <int SmemDpVectorsL>
inline __device__
void commit(
dp_vector_t (&scratch_tile)[BlockDpVectorsK][SmemDpVectorsL])
{
static_assert(SmemDpVectorsL >= BlockDpVectorsL, "Row stride must be >= tile width.");
// Outer thread-tile ldg_vector_t iteration (K-axis)
#pragma unroll
for (int thread_ldgvec_k = 0; thread_ldgvec_k < ThreadLdgVectorsK; ++thread_ldgvec_k)
{
int block_ldgvec_k = block_thread_ldgvec_coords.y + (thread_ldgvec_k * StripmineLdgVectorsK);
// Inner thread-tile ldg_vector_t iteration (L-axis)
#pragma unroll
for (int thread_ldgvec_l = 0; thread_ldgvec_l < ThreadLdgVectorsL; ++thread_ldgvec_l)
{
int block_ldgvec_l = block_thread_ldgvec_coords.x + (thread_ldgvec_l * StripmineLdgVectorsL);
thread_tile[thread_ldgvec_k][thread_ldgvec_l].store(
&scratch_tile[block_ldgvec_k][block_ldgvec_l * LdgVectorDpVectors]);
}
}
}
};
} // namespace gemm
} // namespace cutlass

View File

@ -1,536 +0,0 @@
/******************************************************************************
* Copyright (c) 2011-2017, NVIDIA CORPORATION. All rights reserved.
*
* Redistribution and use in source and binary forms, with or without
* modification, are not permitted.
*
* THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
* ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
* WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
* DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE FOR ANY
* DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
* (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
* LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
* ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
* (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
* SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
*
******************************************************************************/
#pragma once
/**
* \file
* Tile-loading abstraction for thread blocks
*/
#include "../util/util.h"
namespace cutlass {
namespace gemm {
/******************************************************************************
* block_loader (CongruousCopy + idp4 specialization)
******************************************************************************/
/**
* \brief A three-phase data loading abstraction (prefetch, commit, and
* advance) for iterating over ranges of block-wide matrix tiles.
* (CongruousCopy + idp4 specialization)
*
* Each iteration sequence produces a KxL (height-by-width) block-wide tile of
* value_t in shared memory. The layout of the shared block-wide tile is
* a row-major (L-major) tiling of int32_t dp_vector_t, which are themselves
* column-major (K-major) vectors of int8_t value_t. Its dimensions are:
* K = BlockDpVectorsK * (sizeof(dp_vector_t) / sizeof(value_t)
* L = BlockDpVectorsL
*
* The data is copied from a corresponding tile of global matrix data whose
* layout of value_t is also L-major. This constitutes a CongruousCopy between
* the L-major global tile and the L-major shared tile.
*
* NB: The K-major value_t in shared dp_vector_t are imperfectly congruous
* with the L-major value_t in global memory. As a result, the granularity
* of data transfer is a "dp-square" of (DpVectorItems * DpVectorItems) values
* that must be transposed from L-oriented dp_vector_t to K-oriented
* dp_vector_t prior to commitment.
*
* NB: Consecutive threads within a block are mapped in L-major
* fashion across a first-set of squares within their global tile. Successive
* sets of squares are then strip-mined as necessary down the K-axis. These
* discontiguous squares comprise the thread's "slice" of the block-wide tile.
*/
template <
int BlockThreads, ///< Number of threads in each thread block (blockDim.x)
int _BlockDpVectorsK, ///< Extent of block-wide tile in dp_vector_t along the K-axis (height)
int _BlockDpVectorsL, ///< Extent of block-wide tile in dp_vector_t along the L-axis (width)
int LeadingDimAlignBytes, ///< Byte alignment of input matrix leading dimension
bool AllowRaggedTiles ///< Whether the input matrix's dimensions need not be an even-multiple of the block-wide tile dimensions
>
struct block_loader<
BlockThreads,
_BlockDpVectorsK,
_BlockDpVectorsL,
int8_t, ///< Input matrix value type (idp4 specialization)
LeadingDimAlignBytes,
AllowRaggedTiles,
int32_t, ///< Dot-product vector type along the K-axis (idp4 specialization)
load_algorithm::CongruousCopy> ///< Algorithm for loading a shared tile of KxL matrix data (CrosswiseCopy specialization)
{
//-------------------------------------------------------------------------
// Constants and types
//-------------------------------------------------------------------------
/// Input matrix value type
typedef int8_t value_t;
/// Dot-product vector type along the K-axis
typedef int32_t dp_vector_t;
enum
{
/// Number of value_t in a dp_vector_t
DpVectorItems = divide_assert<sizeof(dp_vector_t), sizeof(value_t)>::value,
/// Number of dp_vector_t in a block-wide tile
BlockDpVectors = _BlockDpVectorsK * _BlockDpVectorsL,
/// Number of dp_vector_t in a thread-tile
ThreadDpVectors = divide_assert<BlockDpVectors, BlockThreads>::value,
/// Number of dp_vector_t in a dp-square
SquareDpVectors = DpVectorItems,
/// Number of dp-square tiles in a thread-tile
ThreadSquares = divide_assert<ThreadDpVectors, SquareDpVectors>::value,
/// Extent of block-wide tile in transposed dp_vector_t along the K-axis (height)
BlockTransDpVectorsK = _BlockDpVectorsK * DpVectorItems,
/// Extent of block-wide tile in transposed dp_vector_t along the L-axis (height)
BlockTransDpVectorsL = divide_assert<_BlockDpVectorsL, DpVectorItems>::value,
};
/// Load-from-global data movement type, coarsened by LeadingDimAlignBytes, capped by the
/// smaller of either ThreadSquares or BlockTransDpVectorsL
typedef io_vector<
dp_vector_t,
__NV_STD_MIN(ThreadSquares, BlockTransDpVectorsL),
LeadingDimAlignBytes>
ldg_vector_t;
/// Store-to-shared data movement type equivalent to a dp-square
typedef io_vector<
dp_vector_t,
SquareDpVectors>
sts_vector_t;
enum
{
/// Number of dp_vector_t per ldg_vector_t
LdgVectorDpVectors = ldg_vector_t::VectorItems,
/// Number of value_t per ldg_vector_t
LdgVectorItems = LdgVectorDpVectors * DpVectorItems,
/// Total number of ldg_vector_t within each block-wide tile
BlockLdgVectors = divide_assert<BlockDpVectors, LdgVectorDpVectors>::value,
/// Extent of the block-wide tile in ldg_vector_t along L-axis
BlockLdgVectorsL = divide_assert<BlockTransDpVectorsL, LdgVectorDpVectors>::value,
/// Extent of the block-wide tile in ldg_vector_t along K-axis
BlockLdgVectorsK = BlockTransDpVectorsK,
/// Number of ldg_vector_t within each thread-tile
ThreadLdgVectors = divide_assert<BlockLdgVectors, BlockThreads>::value,
/// Extent of the thread tile in ldg_vector_t along L-axis
ThreadLdgVectorsL = __NV_STD_MAX(1, (BlockLdgVectorsL / BlockThreads)),
/// Extent of the thread tile in ldg_vector_t along K-axis
ThreadLdgVectorsK = divide_assert<ThreadLdgVectors, ThreadLdgVectorsL>::value,
/// Extent of the thread tile in dp-square tiles along K-axis
ThreadSquaresK = divide_assert<ThreadLdgVectorsK, SquareDpVectors>::value,
/// Number of ldg_vector_t within each stripmine-tile
StripmineLdgVectors = BlockThreads * SquareDpVectors,
/// Extent of the stripmine tile in ldg_vector_t along L-axis
StripmineLdgVectorsL = __NV_STD_MIN(BlockLdgVectorsL, BlockThreads),
/// Extent of the stripmine tile in ldg_vector_t along K-axis
StripmineLdgVectorsK = divide_assert<StripmineLdgVectors, StripmineLdgVectorsL>::value,
/// Extent of the stripmine tile in dp-square tiles along K-axis
StripmineSquaresK = divide_assert<StripmineLdgVectorsK, SquareDpVectors>::value,
/// Alignment in dp_vector_t along L needed for committing prefetch
AlignmentDpVectorsL = LdgVectorDpVectors,
};
/// Predicate mask type
typedef uint32_t predicate_mask_t;
//-------------------------------------------------------------------------
// Assert assumptions
//-------------------------------------------------------------------------
static_assert((LeadingDimAlignBytes >= 4) && (LeadingDimAlignBytes % 4 == 0),
"Alignment for matrix operands to IGEMM must be a multiple of 4 bytes.");
static_assert(
(ThreadLdgVectors <= sizeof(predicate_mask_t) * 8),
"Predicate mask type does not contain enough bits for encoding load predicates");
//-------------------------------------------------------------------------
// Members
//-------------------------------------------------------------------------
/// Input pointer to matrix in ldg_vector_t
ldg_vector_t *d_matrix_ldgvecs;
/// Extent of the input matrix in ldg_vector_t along the L-axis
int matrix_ldgvecs_l;
/// Thread block's ending ldg_vector_t coordinate (k) within the input matrix (one-past)
int block_end_ldgvec_k;
/// Predicate bits for guarding ldg_vector_t loads within "whole-k" block-wide tiles
predicate_mask_t guard;
/// Predicate bits for guarding ldg_vector_t loads within the final block-wide "residue" tile
predicate_mask_t residue_guard;
/// Iteration span in "whole-k" block-wide tiles
int wholek_tiles_remaining;
/// Distance in ldg_vector_t within pitched-linear memory between successive coordinates along the K-axis
int matrix_ldgvec_stride_k;
/// Distance in ldg_vector_t within pitched-linear memory between successive coordinates along the L-axis
int matrix_ldgvec_stride_l;
/// ldg_vector_t coordinates (l, k) of thread-tile within the block-wide tile
int2 block_thread_ldgvec_coords;
/// Thread-wide tile of prefetch data
ldg_vector_t thread_tile[ThreadSquaresK][SquareDpVectors][ThreadLdgVectorsL];
//-------------------------------------------------------------------------
// Utility methods
//-------------------------------------------------------------------------
/**
* \brief Byte-permute. Pick four arbitrary bytes from two 32-bit registers, and reassemble them into a 32-bit destination register. For SM2.0 or later.
*
* \par
* The bytes in the two source registers \p a and \p b are numbered from 0 to 7:
* {\p b, \p a} = {{b7, b6, b5, b4}, {b3, b2, b1, b0}}. For each of the four bytes
* {b3, b2, b1, b0} selected in the return value, a 4-bit selector is defined within
* the four lower "nibbles" of \p index: {\p index } = {n7, n6, n5, n4, n3, n2, n1, n0}
*
* \par Snippet
* The code snippet below illustrates byte-permute.
* \par
* \code
* #include <cub/cub.cuh>
*
* __global__ void ExampleKernel(...)
* {
* int a = 0x03020100;
* int b = 0x07060504;
* int index = 0x00007531;
*
* int selected = prmt(a, b, index); // 0x07050301
*
* \endcode
*
*/
inline __device__
int32_t prmt(int32_t a, int32_t b, unsigned int index)
{
int ret;
asm volatile("prmt.b32 %0, %1, %2, %3;" : "=r"(ret) : "r"(a), "r"(b), "r"(index));
return ret;
}
/**
* Convert a "dp-square" from L-major to K-major
*/
inline __device__
void transpose_dp_square(dp_vector_t (&dp_square)[SquareDpVectors])
{
// Transpose dp_vector_t squares
int32_t y = prmt(dp_square[0], dp_square[1], 0x00007362);
int32_t w = prmt(dp_square[2], dp_square[3], 0x00007362);
int32_t x = prmt(dp_square[0], dp_square[1], 0x00005140);
int32_t z = prmt(dp_square[2], dp_square[3], 0x00005140);
dp_square[0] = prmt(x, z, 0x00005410);
dp_square[1] = prmt(x, z, 0x00007632);
dp_square[2] = prmt(y, w, 0x00005410);
dp_square[3] = prmt(y, w, 0x00007632);
}
//-------------------------------------------------------------------------
// Constructor API
//-------------------------------------------------------------------------
/// Constructor
inline __device__
block_loader(
value_t *d_matrix_items, ///< Input pointer to matrix in value_t
int matrix_items_l, ///< Extent of the input matrix in value_t along the L-axis
int matrix_items_stride_k, ///< Distance in value_t within pitched-linear memory between successive coordinates along the K-axis
int matrix_items_stride_l, ///< Distance in value_t within pitched-linear memory between successive coordinates along the L-axis
int2 matrix_block_item_coords, ///< value_t coordinates (l, k) of first block-wide tile within the input matrix
int block_end_item_k) ///< Thread block's ending coordinate (k) within the input matrix (one-past)
:
block_end_ldgvec_k(block_end_item_k),
guard(0),
residue_guard(0)
{
matrix_ldgvecs_l = matrix_items_l / LdgVectorItems;
matrix_ldgvec_stride_k = matrix_items_stride_k / LdgVectorItems,
matrix_ldgvec_stride_l = matrix_items_stride_l;
// ldg_vector_t coordinates (l, k) of thread-tile within the block-wide tile
block_thread_ldgvec_coords = make_int2(
threadIdx.x % BlockLdgVectorsL, // l-coordinate
(threadIdx.x / BlockLdgVectorsL) * SquareDpVectors); // k-coordinate
// ldg_vector_t coordinates (l, k) of first block-wide tile within the input matrix
int2 matrix_block_ldgvec_coords = make_int2(
matrix_block_item_coords.x / LdgVectorItems, // l-coordinate
matrix_block_item_coords.y); // k-coordinate
// Iteration span in ldg_vector_t
int span_ldgvec_k = (block_end_item_k - matrix_block_item_coords.y);
// ldg_vector_t coordinates (l, k) of first thread-tile tile within the input matrix
int2 matrix_thread_ldgvec_coords = make_int2(
block_thread_ldgvec_coords.x + matrix_block_ldgvec_coords.x,
block_thread_ldgvec_coords.y + matrix_block_ldgvec_coords.y);
// Iteration range in "whole-k" block-wide tiles
wholek_tiles_remaining = span_ldgvec_k / BlockLdgVectorsK;
// Extent of final residue-tile in ldg_vector_t along K-axis
int residue_ldgvecs_k = span_ldgvec_k % BlockLdgVectorsK;
// Initialize I/O predicates
if (AllowRaggedTiles)
{
// Iterate through rows of squares in thread tile
#pragma unroll
for (int thread_square_k = 0; thread_square_k < ThreadSquaresK; ++thread_square_k)
{
// Iterate through rows of dp_vector_t in each square
#pragma unroll
for (int square_dpvec = 0; square_dpvec < SquareDpVectors; ++square_dpvec)
{
// ldg_vector_t K-coordinate in block-wide tile (K-axis strip-mining of ldg_vector_t within block-tile)
int block_ldgvec_k =
block_thread_ldgvec_coords.y +
(thread_square_k * StripmineLdgVectorsK) +
square_dpvec;
// Whether block_ldgvec_coords.y is valid in the final residue tile
predicate_mask_t valid_k = (block_ldgvec_k < residue_ldgvecs_k);
// L-axis strip-mining of block-tile
#pragma unroll
for (int thread_ldgvec_l = 0; thread_ldgvec_l < ThreadLdgVectorsL; ++thread_ldgvec_l)
{
// ldg_vector_t L-coordinate in block-wide tile (L-axis strip-mining of ldg_vector_t within block-tile)
int block_ldgvec_l = block_thread_ldgvec_coords.x + (thread_ldgvec_l * StripmineLdgVectorsL);
// Whether block_ldgvec_coords.x is valid any block-wide tile
predicate_mask_t valid_l = (matrix_block_ldgvec_coords.x + block_ldgvec_l < matrix_ldgvecs_l);
// Linear index of ldg_vector_t load
int ldgvec_idx =
(thread_square_k * SquareDpVectors * ThreadLdgVectorsL) +
(square_dpvec * ThreadLdgVectorsL) +
thread_ldgvec_l;
// Set predicate guard bits
guard |= (valid_l << ldgvec_idx);
residue_guard |= ((valid_l & valid_k) << ldgvec_idx);
}
}
}
// Promote residue-guard to primary-guard if no full tiles remain
if (!wholek_tiles_remaining)
{
guard = residue_guard;
}
}
// Update the input pointer to be matrix_thread_ldgvec_coords
this->d_matrix_ldgvecs =
reinterpret_cast<ldg_vector_t*>(d_matrix_items) +
(matrix_thread_ldgvec_coords.y * matrix_ldgvec_stride_k) +
(matrix_thread_ldgvec_coords.x * matrix_ldgvec_stride_l);
}
//-------------------------------------------------------------------------
// Loader API
//-------------------------------------------------------------------------
/**
* Request the current block-wide tile
*/
inline __device__
void request()
{
// Each thread iterates through the ldg_vector_t in its thread tile
// Iterate through rows of squares in thread tile
#pragma unroll
for (int thread_square_k = 0; thread_square_k < ThreadSquaresK; ++thread_square_k)
{
// Iterate through rows of dp_vector_t in each square
#pragma unroll
for (int square_dpvec = 0; square_dpvec < SquareDpVectors; ++square_dpvec)
{
// Iterate through ldg_vector_t in each row
#pragma unroll
for (int thread_ldgvec_l = 0; thread_ldgvec_l < ThreadLdgVectorsL; ++thread_ldgvec_l)
{
// Linear index of ldg_vector_t load
int ldgvec_idx =
(thread_square_k * SquareDpVectors * ThreadLdgVectorsL) +
(square_dpvec * ThreadLdgVectorsL) +
thread_ldgvec_l;
// Unpack predicate guard
predicate_mask_t valid = ((guard >> ldgvec_idx) & 1);
if (!AllowRaggedTiles || valid)
{
// Perform load
thread_tile[thread_square_k][square_dpvec][thread_ldgvec_l].load(
d_matrix_ldgvecs +
(((thread_square_k * StripmineLdgVectorsK) + square_dpvec) * matrix_ldgvec_stride_k) +
(thread_ldgvec_l * StripmineLdgVectorsL * matrix_ldgvec_stride_l));
}
else
{
// Zero-initialize
#pragma unroll
for (int dpvec = 0; dpvec < LdgVectorDpVectors; ++dpvec)
thread_tile[thread_square_k][square_dpvec][thread_ldgvec_l].buff[dpvec] = 0;
}
}
}
}
}
/**
* Advance the loader to the next block-wide tile in the K-axis
*/
inline __device__
void next()
{
d_matrix_ldgvecs += (matrix_ldgvec_stride_k * BlockLdgVectorsK);
if (AllowRaggedTiles)
{
--wholek_tiles_remaining;
// Promote residue-guard to primary-guard if no full tiles remain
if (!wholek_tiles_remaining)
{
guard = residue_guard;
}
}
}
/**
* Commit the previously-requested block-wide tile to shared memory
*
* NB: To facilitate padding for avoiding shared memory bank conflicts, we
* allow the row stride SmemDpVectorsL to be arbitrarily bigger than the
* tile width BlockDpVectorsL.
*/
template <int SmemDpVectorsL>
inline __device__
void commit(
dp_vector_t (&scratch_tile)[_BlockDpVectorsK][SmemDpVectorsL])
{
static_assert(SmemDpVectorsL >= _BlockDpVectorsL, "Row stride must be >= tile width.");
// Square K-coordinate of thread tile in block-wide tile
int block_thread_square_k = block_thread_ldgvec_coords.y / SquareDpVectors;
// Iterate through rows of squares in thread tile
#pragma unroll
for (int thread_square_k = 0; thread_square_k < ThreadSquaresK; ++thread_square_k)
{
// Square K-coordinate in block-wide tile (K-axis strip-mining of squares within block-tile)
int block_square_k = block_thread_square_k + (thread_square_k * StripmineSquaresK);
// Iterate through ldg_vector_t in each row
#pragma unroll
for (int thread_ldgvec_l = 0; thread_ldgvec_l < ThreadLdgVectorsL; ++thread_ldgvec_l)
{
// ldg_vector_t L-coordinate in block-wide tile (L-axis strip-mining of ldg_vector_t within block-tile)
int block_ldgvec_l = block_thread_ldgvec_coords.x + (thread_ldgvec_l * StripmineLdgVectorsL);
// Iterate through squares in each ldg_vector_t
#pragma unroll
for (int ldgvec_dpvec_l = 0; ldgvec_dpvec_l < LdgVectorDpVectors; ++ldgvec_dpvec_l)
{
// Square L-coordinate in block-wide tile (L-axis raking of square-slices within ldg_vector_t)
int block_square_l = (block_ldgvec_l * LdgVectorDpVectors) + ldgvec_dpvec_l;
// Assemble square of L-major dp_vector_t from stack of slices
sts_vector_t square;
// Iterate through rows of dp_vector_t in each square
#pragma unroll
for (int square_dpvec = 0; square_dpvec < SquareDpVectors; ++square_dpvec)
{
square.buff[square_dpvec] = thread_tile[thread_square_k][square_dpvec][thread_ldgvec_l].buff[ldgvec_dpvec_l];
}
// Un-transpose square from L-major to K-major
transpose_dp_square(square.buff);
// Store dp-square
square.store(&scratch_tile[block_square_k][block_square_l * SquareDpVectors]);
}
}
}
}
};
} // namespace gemm
} // namespace cutlass

View File

@ -1,403 +0,0 @@
/******************************************************************************
* Copyright (c) 2011-2017, NVIDIA CORPORATION. All rights reserved.
*
* Redistribution and use in source and binary forms, with or without
* modification, are not permitted.
*
* THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
* ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
* WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
* DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE FOR ANY
* DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
* (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
* LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
* ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
* (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
* SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
*
******************************************************************************/
#pragma once
/**
* \file
* Tile-loading abstraction for thread blocks
*/
#include "../util/util.h"
namespace cutlass {
namespace gemm {
/******************************************************************************
* block_loader (CrosswiseCopy specialization)
******************************************************************************/
/**
* \brief A three-phase data loading abstraction (prefetch, commit, and
* advance) for iterating over ranges of block-wide matrix tiles.
* (CrosswiseCopy specialization)
*
* Each iteration sequence produces a KxL (height-by-width) block-wide tile of
* value_t in shared memory. The layout of the shared block-wide tile is
* a row-major (L-major) tiling of dp_vector_t items, which are themselves
* column-major (K-major) vectors of value_t. Its dimensions are:
* K = BlockDpVectorsK * (sizeof(dp_vector_t) / sizeof(value_t)
* L = BlockDpVectorsL
*
* The data is copied from a corresponding tile of global matrix data whose
* layout of value_t is K-major. This constitutes a CrosswiseCopy between
* the K-major global tile and the L-major shared tile.
*
* NB: The orientation of dp_vector_t components in shared memory is congruous
* with the global matrix data, so we can use dp_vector_t as the minimum
* granularity of data transfer without any intermediate {dis|re}assembly
* of its value_t components. However, the global and shared memory layouts
* of dp_vector_t items are cross-wise with respect to each other, so any
* further LDG-vectorization of dp_vector_t data requires intermediate
* disassembly into dp_vector_t components to be stored individually into
* the shared tile.
*
* NB: Consecutive threads within a block are mapped in K-major
* fashion down a first set of LDG-vectors of dp_vector_t within their global
* tile. Successive sets of LDG-vectors are then strip-mined as necessary
* across the L-axis. These discontiguous LDG-vectors comprise the thread's
* "slice" of the block-wide tile.
*/
template <
int BlockThreads, ///< Number of threads in each thread block (blockDim.x)
int BlockDpVectorsK, ///< Extent of block-wide tile in dp_vector_t along the K-axis (height)
int BlockDpVectorsL, ///< Extent of block-wide tile in dp_vector_t along the L-axis (width)
typename value_t, ///< Input matrix value type
int LeadingDimAlignBytes, ///< Byte alignment of input matrix leading dimension
bool AllowRaggedTiles, ///< Whether the input matrix's dimensions need not be an even-multiple of the block-wide tile dimensions
typename dp_vector_t> ///< Dot-product vector type along the K-axis
struct block_loader<
BlockThreads,
BlockDpVectorsK,
BlockDpVectorsL,
value_t,
LeadingDimAlignBytes,
AllowRaggedTiles,
dp_vector_t,
load_algorithm::CrosswiseCopy> ///< Algorithm for loading a shared tile of KxL matrix data (CrosswiseCopy specialization)
{
//-------------------------------------------------------------------------
// Constants and types
//-------------------------------------------------------------------------
enum
{
/// Number of value_t in a dp_vector_t
DpVectorItems = divide_assert<sizeof(dp_vector_t), sizeof(value_t)>::value,
/// Number of dp_vector_t in a block-wide tile
BlockDpVectors = BlockDpVectorsK * BlockDpVectorsL,
/// Number of dp_vector_t in a thread-tile
ThreadDpVectors = divide_assert<BlockDpVectors, BlockThreads>::value,
};
/// Data movement type, coarsened by LeadingDimAlignBytes, capped by the
/// smaller of either ThreadDpVectors or BlockDpVectorsK
typedef io_vector<
dp_vector_t,
__NV_STD_MIN(ThreadDpVectors, BlockDpVectorsK),
LeadingDimAlignBytes>
ldg_vector_t;
enum
{
/// Number of dp_vector_t per ldg_vector_t
LdgVectorDpVectors = ldg_vector_t::VectorItems,
/// Number of value_t per ldg_vector_t
LdgVectorItems = LdgVectorDpVectors * DpVectorItems,
/// Total number of ldg_vector_t within each block-wide tile
BlockLdgVectors = divide_assert<BlockDpVectors, LdgVectorDpVectors>::value,
/// Extent of the block-wide tile in ldg_vector_t along K-axis
BlockLdgVectorsK = divide_assert<BlockDpVectorsK, LdgVectorDpVectors>::value,
/// Extent of the block-wide tile in ldg_vector_t along L-axis
BlockLdgVectorsL = BlockDpVectorsL,
/// Number of ldg_vector_t within each thread-tile
ThreadLdgVectors = divide_assert<BlockLdgVectors, BlockThreads>::value,
/// Extent of the thread tile in ldg_vector_t along K-axis
ThreadLdgVectorsK = __NV_STD_MAX(1, (BlockLdgVectorsK / BlockThreads)),
/// Extent of the thread tile in ldg_vector_t along L-axis
ThreadLdgVectorsL = divide_assert<ThreadLdgVectors, ThreadLdgVectorsK>::value,
/// Number of ldg_vector_t within each stripmine-tile
StripmineLdgVectors = BlockThreads,
/// Extent of the stripmine tile in ldg_vector_t along K-axis
StripmineLdgVectorsK = __NV_STD_MIN(BlockLdgVectorsK, StripmineLdgVectors),
/// Extent of the stripmine tile in ldg_vector_t along L-axis
StripmineLdgVectorsL = divide_assert<StripmineLdgVectors, StripmineLdgVectorsK>::value,
/// Alignment in dp_vector_t along L needed for committing prefetch
AlignmentDpVectorsL = 1,
};
/// Predicate bit vector
typedef uint64_t predicate_mask_t;
//-------------------------------------------------------------------------
// Assert assumptions
//-------------------------------------------------------------------------
static_assert(
(ThreadLdgVectors <= sizeof(predicate_mask_t) * 8),
"Predicate mask type does not contain enough bits for encoding load predicates");
//-------------------------------------------------------------------------
// Members
//-------------------------------------------------------------------------
/// Input pointer to matrix in ldg_vector_t
ldg_vector_t *d_matrix_ldgvecs;
/// Extent of the input matrix in ldg_vector_t along the L-axis
int matrix_ldgvecs_l;
/// Thread block's ending ldg_vector_t coordinate (k) within the input matrix (one-past)
int block_end_ldgvec_k;
/// Predicate bits for guarding ldg_vector_t loads within "whole-k" block-wide tiles
predicate_mask_t guard;
/// Predicate bits for guarding ldg_vector_t loads within the final block-wide "residue" tile
predicate_mask_t residue_guard;
/// Iteration span in "whole-k" block-wide tiles
int wholek_tiles_remaining;
/// Distance in ldg_vector_t within pitched-linear memory between successive coordinates along the K-axis
int matrix_ldgvec_stride_k;
/// Distance in ldg_vector_t within pitched-linear memory between successive coordinates along the L-axis
int matrix_ldgvec_stride_l;
/// ldg_vector_t coordinates (l, k) of thread-tile within the block-wide tile
int2 block_thread_ldgvec_coords;
/// Thread-wide tile of prefetch data
ldg_vector_t thread_tile[ThreadLdgVectorsK][ThreadLdgVectorsL];
//-------------------------------------------------------------------------
// Constructor API
//-------------------------------------------------------------------------
/// Constructor
inline __device__
block_loader(
value_t *d_matrix_items, ///< Input pointer to matrix in value_t
int matrix_items_l, ///< Extent of the input matrix in value_t along the L-axis
int matrix_items_stride_k, ///< Distance in value_t within pitched-linear memory between successive coordinates along the K-axis
int matrix_items_stride_l, ///< Distance in value_t within pitched-linear memory between successive coordinates along the L-axis
int2 matrix_block_item_coords, ///< value_t coordinates (l, k) of first block-wide tile within the input matrix
int block_end_item_k) ///< Thread block's ending coordinate (k) within the input matrix (one-past)
:
block_end_ldgvec_k(block_end_item_k),
guard(0),
residue_guard(0)
{
matrix_ldgvecs_l = matrix_items_l;
matrix_ldgvec_stride_k = matrix_items_stride_k;
matrix_ldgvec_stride_l = (matrix_items_stride_l / LdgVectorItems);
// ldg_vector_t coordinates (l, k) of thread-tile within the block-wide tile
block_thread_ldgvec_coords = make_int2(
(threadIdx.x / BlockLdgVectorsK), // l-coordinate
(threadIdx.x % BlockLdgVectorsK)); // k-coordinate
// ldg_vector_t coordinates (l, k) of first block-wide tile within the input matrix
int2 matrix_block_ldgvec_coords = make_int2(
matrix_block_item_coords.x, // l-coordinate
matrix_block_item_coords.y / LdgVectorItems); // k-coordinate
// Iteration span in ldg_vector_t
int span_ldgvec_k = (block_end_item_k - matrix_block_item_coords.y) / LdgVectorItems;
// ldg_vector_t coordinates (l, k) of first thread-tile tile within the input matrix
int2 matrix_thread_ldgvec_coords = make_int2(
block_thread_ldgvec_coords.x + matrix_block_ldgvec_coords.x,
block_thread_ldgvec_coords.y + matrix_block_ldgvec_coords.y);
// Iteration range in "whole-k" block-wide tiles
wholek_tiles_remaining = span_ldgvec_k / BlockLdgVectorsK;
// Extent of final residue-tile in ldg_vector_t along K-axis
int residue_ldgvecs_k = span_ldgvec_k % BlockLdgVectorsK;
// Initialize I/O predicates
if (AllowRaggedTiles)
{
// Outer thread-tile ldg_vector_t iteration (K-axis)
#pragma unroll
for (int thread_ldgvec_k = 0; thread_ldgvec_k < ThreadLdgVectorsK; ++thread_ldgvec_k)
{
int block_ldgvec_k = block_thread_ldgvec_coords.y + (thread_ldgvec_k * StripmineLdgVectorsK);
// Whether block_ldgvec_coords.y is valid in the final residue tile
predicate_mask_t valid_k = (block_ldgvec_k < residue_ldgvecs_k);
// Inner thread-tile ldg_vector_t iteration (L-axis)
#pragma unroll
for (int thread_ldgvec_l = 0; thread_ldgvec_l < ThreadLdgVectorsL; ++thread_ldgvec_l)
{
int block_ldgvec_l = block_thread_ldgvec_coords.x + (thread_ldgvec_l * StripmineLdgVectorsL);
// Whether block_ldgvec_coords.x is valid any block-wide tile
predicate_mask_t valid_l = (matrix_block_ldgvec_coords.x + block_ldgvec_l < matrix_ldgvecs_l);
// Linear index of ldg_vector_t load
int ldgvec_idx = thread_ldgvec_l + (thread_ldgvec_k * ThreadLdgVectorsL);
// Set predicate guard bits
guard |= (valid_l << ldgvec_idx);
residue_guard |= ((valid_l & valid_k) << ldgvec_idx);
}
}
// Promote residue-guard to primary-guard if no full tiles remain
if (!wholek_tiles_remaining)
{
guard = residue_guard;
}
}
// Update the input pointer to be matrix_thread_ldgvec_coords
this->d_matrix_ldgvecs =
reinterpret_cast<ldg_vector_t*>(d_matrix_items) +
(matrix_thread_ldgvec_coords.y * matrix_ldgvec_stride_k) +
(matrix_thread_ldgvec_coords.x * matrix_ldgvec_stride_l);
}
//-------------------------------------------------------------------------
// Loader API
//-------------------------------------------------------------------------
/**
* Request the current block-wide tile
*/
inline __device__
void request()
{
// Outer thread-tile ldg_vector_t iteration (K-axis)
#pragma unroll
for (int thread_ldgvec_k = 0; thread_ldgvec_k < ThreadLdgVectorsK; ++thread_ldgvec_k)
{
// Inner thread-tile ldg_vector_t iteration (L-axis)
#pragma unroll
for (int thread_ldgvec_l = 0; thread_ldgvec_l < ThreadLdgVectorsL; ++thread_ldgvec_l)
{
// Linear index of ldg_vector_t load
int ldgvec_idx = (thread_ldgvec_k * ThreadLdgVectorsL) + thread_ldgvec_l;
// Unpack predicate guard
predicate_mask_t valid = ((guard >> ldgvec_idx) & 1);
if (!AllowRaggedTiles || valid)
{
// Perform load
thread_tile[thread_ldgvec_k][thread_ldgvec_l].load(
d_matrix_ldgvecs +
(thread_ldgvec_k * StripmineLdgVectorsK * matrix_ldgvec_stride_k) +
(thread_ldgvec_l * StripmineLdgVectorsL * matrix_ldgvec_stride_l));
}
else
{
// Zero-initialize
#pragma unroll
for (int dpvec = 0; dpvec < LdgVectorDpVectors; ++dpvec)
thread_tile[thread_ldgvec_k][thread_ldgvec_l].buff[dpvec] = 0;
}
}
}
}
/**
* Advance the loader to the next block-wide tile in the K-axis
*/
inline __device__
void next()
{
d_matrix_ldgvecs += (matrix_ldgvec_stride_k * BlockLdgVectorsK);
if (AllowRaggedTiles)
{
--wholek_tiles_remaining;
// Promote residue-guard to primary-guard if no full tiles remain
if (!wholek_tiles_remaining)
{
guard = residue_guard;
}
}
}
/**
* Commit the previously-requested block-wide tile to shared memory
*
* NB: To facilitate padding for avoiding shared memory bank conflicts, we
* allow the row stride SmemDpVectorsL to be arbitrarily bigger than the
* tile width BlockDpVectorsL.
*/
template <int SmemDpVectorsL>
inline __device__
void commit(
dp_vector_t (&scratch_tile)[BlockDpVectorsK][SmemDpVectorsL])
{
static_assert(SmemDpVectorsL >= BlockDpVectorsL, "Row stride must be >= tile width.");
// Outer thread-tile ldg_vector_t iteration (K-axis)
#pragma unroll
for (int thread_ldgvec_k = 0; thread_ldgvec_k < ThreadLdgVectorsK; ++thread_ldgvec_k)
{
int block_ldgvec_k = block_thread_ldgvec_coords.y + (thread_ldgvec_k * StripmineLdgVectorsK);
// Inner thread-tile ldg_vector_t iteration (L-axis)
#pragma unroll
for (int thread_ldgvec_l = 0; thread_ldgvec_l < ThreadLdgVectorsL; ++thread_ldgvec_l)
{
int block_ldgvec_l = block_thread_ldgvec_coords.x + (thread_ldgvec_l * StripmineLdgVectorsL);
// Write column of dp_vector_t
#pragma unroll
for (int dpvec = 0; dpvec < LdgVectorDpVectors; ++dpvec)
{
scratch_tile[(block_ldgvec_k * LdgVectorDpVectors) + dpvec][block_ldgvec_l] =
thread_tile[thread_ldgvec_k][thread_ldgvec_l].buff[dpvec];
}
}
}
}
};
} // namespace gemm
} // namespace cutlass

View File

@ -1,314 +0,0 @@
/******************************************************************************
* Copyright (c) 2011-2017, NVIDIA CORPORATION. All rights reserved.
*
* Redistribution and use in source and binary forms, with or without
* modification, are not permitted.
*
* THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
* ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
* WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
* DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE FOR ANY
* DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
* (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
* LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
* ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
* (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
* SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
*
******************************************************************************/
#pragma once
/**
* \file
* Tile-loading abstraction for thread blocks
*/
#include "../util/util.h"
namespace cutlass {
namespace gemm {
/**
* block-wide tile loader supporting congruous mapping of data from source and
* destination addressable storage. Typically, this will be used to load a
* block-wide tile from global memory into shared memory.
*
* This enables the caller to specify MatrixAlignBytes guarantees of the input pointer
* and performs memory operations on vectors. This increases the efficiency of
* memory operations and reduces the number of guard predicates needed.
*
*/
template <
bool congruous, ///< Indicates whether the "GEMM K" dimension refers to strided matrix dimension
int BlockThreads, ///< Number of threads participating in the streaming operation
int BlockItemsL, ///< Extent of block-wide tile in value_t along the L-axis (width)
int BlockItemsK, ///< Extent of block-wide tile in value_t along the K-axis (height)
typename value_t, ///< Input matrix value type
int MatrixAlignBytes, ///< Byte alignment of input matrix
bool AllowRaggedTiles ///< Whether the input matrix's dimensions need not be an even-multiple of the block-wide tile dimensions
>
struct block_loader_wmma
{
//-------------------------------------------------------------------------
// Constants and types
//-------------------------------------------------------------------------
/// Predicate bit vector
typedef uint64_t predicate_mask_t;
/// Data movement type, coarsened by MatrixAlignBytes
typedef io_vector<
value_t,
divide_assert<MatrixAlignBytes, sizeof(value_t)>::value,
MatrixAlignBytes>
ldg_vector_t;
enum
{
/// Number of items per ldg_vector_t
LdgVectorItems = ldg_vector_t::VectorItems,
/// Total number of ldg_vector_t within the block-wide tile
BlockLdgVectors = divide_assert<(BlockItemsL * BlockItemsK), LdgVectorItems>::value,
/// Extent of the block-wide tile in ldg_vector_t along K-axis
BlockLdgVectorsK = BlockItemsK,
/// Extent of the block-wide tile in ldg_vector_t along L-axis
BlockLdgVectorsL = divide_assert<BlockItemsL, LdgVectorItems>::value,
/// Number of ldg_vector_t within each thread tile
ThreadLdgVectors = divide_assert<BlockLdgVectors, BlockThreads>::value,
/// Extent of the thread tile in ldg_vector_t along the L-axis
ThreadLdgVectorsL = __NV_STD_MAX(1, BlockLdgVectorsL / BlockThreads),
/// Block-wide strip-mining distance between ldg_vector_t along the K-axis
BlockLdgVectorStrideK = __NV_STD_MAX(1, BlockThreads / BlockLdgVectorsL),
/// Extent of the thread tile in ldg_vector_t along the K-axis
ThreadLdgVectorsK = divide_assert<BlockLdgVectorsK, BlockLdgVectorStrideK>::value,
};
//-------------------------------------------------------------------------
// Assert assumptions
//-------------------------------------------------------------------------
/// Define assertions
static_assert(ThreadLdgVectorsL * ThreadLdgVectorsK == ThreadLdgVectors,
"Number of vectors must be fully covered by the thread's 2D vector tile.");
/// Predicate masks must be large enough to guard every vector load
static_assert(sizeof(predicate_mask_t) * 8 >= ThreadLdgVectorsL * ThreadLdgVectorsK,
"Predicate bit vector must be large enough to guard every vector load.");
//-------------------------------------------------------------------------
// Members
//-------------------------------------------------------------------------
/// pointer to tile in global memory
const ldg_vector_t *ptr;
/// stride of the matrix in the K-axis
int matrix_values_stride_k;
/// Guard predicate
predicate_mask_t guard;
/// Guard for the last request iteration
predicate_mask_t residue_guard;
/// Number of 'whole' request iterations before encountering the residue
int request_iterations;
/// fetch registers
ldg_vector_t fetch[ThreadLdgVectors];
/// Thread's base offset from the start of a block-wide tile
int thread_offset_l;
/// Thread's basae offset from the start of a block-wide tile
int thread_offset_k;
//-------------------------------------------------------------------------
// Constructor API
//-------------------------------------------------------------------------
/// Constructor
inline __device__
block_loader_wmma(
const value_t *d_matrix, ///< Pointer to input matrix
int matrix_values_l, ///< Extent of the input matrix in value_t along the L-axis
int start_l, ///< Starting location in tile
int dim_k, ///< Inner dimension of tile, used for computing guard predicates
int _matrix_values_stride_k, ///< Stride of K-axis of atrix
int start_k, ///< Tile's starting location
int2 block_begin_item_coords) ///< Thread block's starting value_t coordinates (l, k) within the input matrix
:
ptr(reinterpret_cast<const ldg_vector_t *>(d_matrix)),
matrix_values_stride_k(_matrix_values_stride_k / LdgVectorItems),
guard(0),
residue_guard(0)
{
// Compute block's starting coordinates in units of vectors
int block_base_l = block_begin_item_coords.x / LdgVectorItems;
int block_base_k = block_begin_item_coords.y;
// Compute a thread tiling of the block-wide tile
int tid = threadIdx.x;
thread_offset_l = tid % BlockLdgVectorsL;
thread_offset_k = tid / BlockLdgVectorsL;
// Add the block and thread offsets to the source pointer
ptr += (block_base_l + thread_offset_l) +
(block_base_k + thread_offset_k) * matrix_values_stride_k;
// When AllowRaggedTiles support is enabled, compute a bit vector of guard
// predicates
if (AllowRaggedTiles)
{
if (congruous)
{
request_iterations = (dim_k - start_k) / BlockItemsK;
}
else
{
request_iterations = (matrix_values_l - start_l) / BlockItemsL;
}
#pragma unroll
for (int k_idx = 0; k_idx < ThreadLdgVectorsK; ++k_idx)
{
#pragma unroll
for (int l_idx = 0; l_idx < ThreadLdgVectorsL; ++l_idx)
{
int item = l_idx + k_idx * ThreadLdgVectorsL;
// Global vector L and K indices
int vec_l = l_idx * BlockThreads;
int vec_k = k_idx * BlockLdgVectorStrideK;
predicate_mask_t pred;
predicate_mask_t residue_pred;
if (congruous)
{
pred = (((block_base_l + thread_offset_l + vec_l) * LdgVectorItems < matrix_values_l) ? 1 : 0);
residue_pred = ((block_base_k + thread_offset_k + vec_k < (dim_k % BlockItemsK)) ? 1 : 0);
}
else
{
pred = ((block_base_k + thread_offset_k + vec_k < dim_k) ? 1 : 0);
residue_pred = (((block_base_l + thread_offset_l + vec_l) * LdgVectorItems < (matrix_values_l % BlockItemsL)) ? 1 : 0);
}
// Update the guard and residue_guard word with predicate bits
guard |= (pred << item);
residue_guard |= (residue_pred << item);
}
}
// If there are zero full request iterations, compute the intersection
// with the residue guard.
if (!request_iterations)
{
guard &= residue_guard;
}
}
}
/**
* Request the current block-wide tile from source memory
*/
inline __device__
void request()
{
#pragma unroll
for (int k_idx = 0; k_idx < ThreadLdgVectorsK; ++k_idx)
{
#pragma unroll
for (int l_idx = 0; l_idx < ThreadLdgVectorsL; ++l_idx)
{
int load_idx = l_idx + (k_idx * ThreadLdgVectorsL);
bool pred = !AllowRaggedTiles || (guard & (predicate_mask_t(1) << load_idx));
if (pred)
{
fetch[load_idx].load(
ptr +
(k_idx * BlockLdgVectorStrideK * matrix_values_stride_k) + (l_idx * BlockThreads));
}
else
{
#pragma unroll
for (int elem_idx = 0; elem_idx < LdgVectorItems; ++elem_idx)
{
fetch[load_idx].buff[elem_idx] = 0;
}
}
}
}
}
/// Advance to the next block-wide tile
inline __device__
void next()
{
if (congruous)
{
ptr += BlockItemsK * matrix_values_stride_k;
}
else
{
ptr += BlockLdgVectorsL;
}
// Track number of full iterations to intersect with the residue guard predicates.
if (AllowRaggedTiles)
{
--request_iterations;
if (!request_iterations)
{
guard &= residue_guard;
}
}
}
/// Commit the values to the scratch tile to destination memory.
template <int SmemStride>
inline __device__
void commit(value_t *scratch_tile)
{
static_assert(SmemStride % LdgVectorItems == 0,
"SMEM stride must be divisible by the size of vector loads");
ldg_vector_t *smem_ptr = reinterpret_cast<ldg_vector_t *>(scratch_tile);
smem_ptr += thread_offset_l + thread_offset_k * SmemStride / LdgVectorItems;
#pragma unroll
for (int k_idx = 0; k_idx < ThreadLdgVectorsK; ++k_idx)
{
#pragma unroll
for (int l_idx = 0; l_idx < ThreadLdgVectorsL; ++l_idx)
{
int load_idx = l_idx + (k_idx * ThreadLdgVectorsL);
fetch[load_idx].store(smem_ptr +
(k_idx * BlockLdgVectorStrideK * SmemStride / LdgVectorItems) +
(l_idx * BlockThreads));
}
}
}
};
} // namespace gemm
} // namespace cutlass

View File

@ -1,669 +0,0 @@
/******************************************************************************
* Copyright (c) 2011-2017, NVIDIA CORPORATION. All rights reserved.
*
* Redistribution and use in source and binary forms, with or without
* modification, are not permitted.
*
* THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
* ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
* WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
* DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE FOR ANY
* DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
* (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
* LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
* ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
* (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
* SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
*
******************************************************************************/
#pragma once
/**
* \file
* A block-wide task abstraction for computing device-wide GEMM
*/
#include <stdint.h>
#include "../util/util.h"
#include "grid_raster.h"
#include "block_loader.h"
#include "k_split_control.h"
#include "thread_accumulator.h"
namespace cutlass {
namespace gemm {
/******************************************************************************
* block_task_policy
******************************************************************************/
/**
* \brief Parameterizable tuning policy for \p block_task
*
* Once parameterized, \p block_task_policy provides the member constant
* \p BlockThreads indicating to the required thread block size
*/
template <
int _BlockItemsY, ///< Height in rows of a block-wide tile in matrix C
int _BlockItemsX, ///< Width in columns of a block-wide tile in matrix C
int _BlockItemsK, ///< Extent of block-wide A|B tiles in value_t along the K-axis
int _ThreadItemsY, ///< Height in rows of a thread tile in C
int _ThreadItemsX, ///< Width in columns of a thread tile in C
bool _UseDoubleScratchTiles, ///< Whether to halve synchronization overhead at the expense of doubled shared memory and addressing overhead
grid_raster_strategy::kind_t _RasterStrategy> ///< Strategy for enumerating \p block_task within an input matrix
struct block_task_policy
{
enum
{
/// Height in rows of a block-wide tile in matrix C
BlockItemsY = _BlockItemsY,
/// Width in columns of a block-wide tile in matrix C
BlockItemsX = _BlockItemsX,
/// Height in rows of a thread tile in C
ThreadItemsY = _ThreadItemsY,
/// Width in columns of a thread tile in C
ThreadItemsX = _ThreadItemsX,
/// Extent of block-wide A|B tiles in value_t along the K-axis
BlockItemsK = _BlockItemsK,
/// Whether to halve synchronization overhead at the expense of doubled shared memory and addressing overhead
UseDoubleScratchTiles = _UseDoubleScratchTiles,
/// Number of threads in each thread block (blockDim.x)
BlockThreads = divide_assert<
(BlockItemsY * BlockItemsX),
(ThreadItemsY * ThreadItemsX)>::value,
};
/// Strategy for enumerating \p block_task within an input matrix
static const grid_raster_strategy::kind_t RasterStrategy = _RasterStrategy;
};
/******************************************************************************
* block_task
******************************************************************************/
/**
* \brief A block-wide task abstraction for computing device-wide GEMM
*
* Each thread_block is assigned a unique tile of output matrix C to compute by
* consuming the corresponding stripes of the input matrices A and B.
*/
template <
typename block_task_policy_t, ///< Parameterization of block_task_policy
typename value_t, ///< Multiplicand value type (matrices A and B)
typename accum_t, ///< Accumulator value type (matrix C and scalars)
matrix_transform_t::kind_t TransformA, ///< View transform enumerant for matrix A
int LdgAlignA, ///< Alignment (in bytes) for A operand
matrix_transform_t::kind_t TransformB, ///< View transform enumerant for matrix B
int LdgAlignB, ///< Alignment (in bytes) for B operand
typename epilogue_op_t, ///< Epilogue operation applied to GEMM
int LdgAlignC, ///< Alignment (in bytes) for C operand
bool AllowRaggedTiles ///< Whether the input matrix's dimensions need not be an even-multiple of the block-wide tile dimensions
>
struct block_task
{
//-------------------------------------------------------------------------
// Constants and types
//-------------------------------------------------------------------------
enum
{
/// Number of threads in each thread block (blockDim.x)
BlockThreads = block_task_policy_t::BlockThreads,
/// Extent of thread tile in value_t along M-axis
ThreadItemsY = block_task_policy_t::ThreadItemsY,
/// Extent of thread tile in value_t along N-axis
ThreadItemsX = block_task_policy_t::ThreadItemsX,
};
/// Accumulator type
typedef thread_accumulator<
ThreadItemsY,
ThreadItemsX,
value_t,
accum_t>
thread_accumulator_t;
/// Dot-product vector type along the K-axis (e.g, uchar4 when using IDP4A)
typedef typename thread_accumulator_t::dp_vector_t dp_vector_t;
enum
{
/// Whether this is a small, latency-bound tile
IsSmallTile = (ThreadItemsY < 4) && (ThreadItemsX < 4),
/// Number of value_t in dp_vector_t
DpVectorItems = divide_assert<sizeof(dp_vector_t), sizeof(value_t)>::value,
/// Extent of block-wide C-tile in accum_t (and A-tiles in value_t) along M-axis (height)
BlockItemsY = block_task_policy_t::BlockItemsY,
/// Extent of block-wide C-tile in accum_t (and B-tiles in value_t) along N-axis (width)
BlockItemsX = block_task_policy_t::BlockItemsX,
/// Extent of block-wide A|B tiles in value_t along the K-axis
BlockItemsK = block_task_policy_t::BlockItemsK,
/// Whether to halve synchronization overhead at the expense of doubled shared memory and addressing overhead
UseDoubleScratchTiles = block_task_policy_t::UseDoubleScratchTiles,
/// Extent of block-wide A|B tiles in dp_vector_t along the K-axis
BlockDpVectorsK = divide_assert<BlockItemsK, DpVectorItems>::value,
/// Number of dp_vector_t along M-axis that can be read in a single LDS from the shared A-tile (up to 128b if more than one value_t)
LdsVectorDpVectorsA = __NV_STD_MIN(
ThreadItemsY,
__NV_STD_MAX(1, (128 / (__NV_STD_MAX(sizeof(dp_vector_t), sizeof(accum_t)) * 8)))),
/// Number of dp_vector_t along N-axis that can be read in a single LDS from the shared B-tile (up to 128b if more than one value_t)
LdsVectorDpVectorsB = __NV_STD_MIN(
ThreadItemsX,
__NV_STD_MAX(1, (128 / (__NV_STD_MAX(sizeof(dp_vector_t), sizeof(accum_t)) * 8)))),
/// Number of strip-mined LDS vector reads from shared A-tile
ThreadLdsVectorsA = divide_assert<ThreadItemsY, LdsVectorDpVectorsA>::value,
/// Number of strip-mined LDS vector reads from shared B-tile
ThreadLdsVectorsB = divide_assert<ThreadItemsX, LdsVectorDpVectorsB>::value,
/// Number of elements in one LDG/STG vector of C-tile
ThreadLdgVectorSizeC = __NV_STD_MIN(LdgAlignC, 16) / (sizeof(accum_t)),
/// Number of threads in warp
WarpThreads = 32,
/// Extent of warp in threads along the M-axis
WarpThreadsY = (BlockItemsY > BlockItemsX) ? 8 : 4,
/// Extent of warp in threads along the N-axis
WarpThreadsX = divide_assert<WarpThreads, WarpThreadsY>::value,
/// Extent of warp-wide tile in items along the M-axis
WarpItemsY = WarpThreadsY * ThreadItemsY,
/// Extent of warp-wide tile in items along the N-axis
WarpItemsX = WarpThreadsX * ThreadItemsX,
/// Extent of block in warps along M-axis
BlockWarpsY = divide_assert<BlockItemsY, WarpItemsY>::value,
/// Extent of block in warps along N-axis
BlockWarpsX = divide_assert<BlockItemsX, WarpItemsX>::value,
};
/// Load-from-shared data movement type for A-tile, coarsened by LdsVectorDpVectorsA
typedef io_vector<dp_vector_t, LdsVectorDpVectorsA> lds_vector_a_t;
/// Load-from-shared data movement type for B-tile, coarsened by LdsVectorDpVectorsB
typedef io_vector<dp_vector_t, LdsVectorDpVectorsB> lds_vector_b_t;
/// Thread block rasterization helper type
typedef grid_raster<
BlockItemsY,
BlockItemsX,
TransformA,
TransformB,
block_task_policy_t::RasterStrategy>
grid_raster_t;
/// Tile loader type for matrix A
typedef block_loader<
BlockThreads, // BlockThreads
BlockDpVectorsK, // BlockDpVectorsK
BlockItemsY, // BlockItemsL
value_t, // value_t
LdgAlignA, // MatrixAlignBytes
AllowRaggedTiles, // AllowRaggedTiles
dp_vector_t, // dp_vector_t
(TransformA == matrix_transform_t::NonTranspose) ? // LoadAlgorithm
load_algorithm::CongruousCopy :
load_algorithm::CrosswiseCopy>
block_loader_a_t;
/// Tile loader type for matrix B
typedef block_loader<
BlockThreads, // BlockThreads
BlockDpVectorsK, // BlockDpVectorsK
BlockItemsX, // BlockItemsL
value_t, // value_t
LdgAlignB, // MatrixAlignBytes
AllowRaggedTiles, // AllowRaggedTiles
dp_vector_t, // dp_vector_t
(TransformB == matrix_transform_t::NonTranspose) ? // LoadAlgorithm
load_algorithm::CrosswiseCopy :
load_algorithm::CongruousCopy>
block_loader_b_t;
enum
{
/// Number of value_t to pad the end of each row of the shared A-tile
PadItemsA = (TransformA == matrix_transform_t::NonTranspose) ?
__NV_STD_MAX(LdsVectorDpVectorsA, block_loader_a_t::AlignmentDpVectorsL) :
LdsVectorDpVectorsA,
/// Number of value_t to pad the end of each row of the shared B-tile
PadItemsB = (TransformB == matrix_transform_t::NonTranspose) ?
LdsVectorDpVectorsB :
__NV_STD_MAX(LdsVectorDpVectorsB, block_loader_b_t::AlignmentDpVectorsL),
};
/// Shared memory layout for a prefetch page
struct page_storage_t
{
/// Tile of A
dp_vector_t __align__(16) block_a[BlockDpVectorsK][BlockItemsY + PadItemsA];
/// Tile of B
dp_vector_t __align__(16) block_b[BlockDpVectorsK][BlockItemsX + PadItemsB];
};
/// Shared memory layout for scratch storage
struct scratch_storage_t
{
/// Prefetch pages
page_storage_t pages[UseDoubleScratchTiles ? 2 : 1];
/// Accumulator shared scratch
typename thread_accumulator_t::scratch_storage_t accum_scratch;
};
//-------------------------------------------------------------------------
// Assert assumptions
//-------------------------------------------------------------------------
// Ensure we have at least two unrolled innermost loop iterations (one to prefetch
// the next global tile and then one to prefetch the first strip of it from shared)
static_assert ((BlockDpVectorsK >= 2), "BlockDpVectorsK must be >= 2.");
//-------------------------------------------------------------------------
// Members
//-------------------------------------------------------------------------
/// Scratch storage reference
scratch_storage_t *scratch;
/// Which page of scratch tiles we're currently reading from
int page_idx;
/// Pointer to matrix C
accum_t *d_c;
/// Epilogue operation applied to update matrix C
epilogue_op_t epilogue_op;
/// Matrix height in rows of trans_op(A) and C
int dim_m;
/// Matrix width in columns of trans_op(B) and C
int dim_n;
/// Control for inter-block k-splitting
k_split_control k_split;
/// Thread block's base value_t coordinates (m, n) in matrix C
grid_raster_t grid_raster;
/// Thread block's current coordinate (k) within A|B matrices
int block_item_coords_k;
/// Thread block's ending coordinate (k) within A|B matrices (one-past)
int block_end_item_k;
/// Warp's coordinates (x, y) in thread block
int2 block_warp_coords;
/// Thread's coordinates (x, y) in warp
int2 warp_thread_coords;
/// Thread's base item offset within strip of A tile
int thread_strip_offset_a;
/// Thread's base item offset within strip of B tile
int thread_strip_offset_b;
/// Thread's active-k/prefetch-k slices from shared A tile
lds_vector_a_t local_slices_a[2][ThreadLdsVectorsA];
/// Thread's active-k/prefetch-k slices from shared B tile
lds_vector_b_t local_slices_b[2][ThreadLdsVectorsB];
/// A tile loader
block_loader_a_t loader_a;
/// B tile loader
block_loader_b_t loader_b;
/// C tile accumulator
thread_accumulator_t accumulator;
//-------------------------------------------------------------------------
// Coordinate system helpers
//-------------------------------------------------------------------------
/// Compute the warp's coordinates (x, y) in thread block
inline __device__
int2 warp_coords()
{
int warp_id = threadIdx.x / WarpThreads;
return make_int2(
warp_id % BlockWarpsX,
warp_id / BlockWarpsX);
}
/// Compute the thread's lane-coordinates (x, y) in warp
inline __device__
int2 thread_coords()
{
int lane_id = threadIdx.x % WarpThreads;
// Maxwell+ mapping of threads within a 2D warp for maximal LDS bandwidth
return make_int2(
lane_id / WarpThreadsY,
lane_id % WarpThreadsY);
}
//-------------------------------------------------------------------------
// Constructor API
//-------------------------------------------------------------------------
/// Constructor
inline __device__
block_task(
scratch_storage_t *scratch,
value_t *d_a,
value_t *d_b,
accum_t *d_c,
epilogue_op_t epilogue_op,
int dim_m,
int dim_n,
int dim_k,
k_split_control k_split)
:
scratch(scratch),
page_idx(0),
d_c(d_c),
epilogue_op(epilogue_op),
dim_m(dim_m),
dim_n(dim_n),
k_split(k_split),
block_item_coords_k(k_split.block_begin_item_k()),
block_end_item_k(k_split.block_end_item_k(dim_k)),
block_warp_coords(warp_coords()),
warp_thread_coords(thread_coords()),
thread_strip_offset_a((warp_thread_coords.y * LdsVectorDpVectorsA) + (block_warp_coords.y * WarpItemsY)),
thread_strip_offset_b((warp_thread_coords.x * LdsVectorDpVectorsB) + (block_warp_coords.x * WarpItemsX)),
loader_a(
d_a, // d_matrix
dim_m, // matrix_values_l
(TransformA == matrix_transform_t::NonTranspose) ? dim_m : 1, // matrix_values_stride_k
(TransformA == matrix_transform_t::NonTranspose) ? 1 : dim_k, // matrix_values_stride_l
make_int2( // block_begin_item_coords
grid_raster.block_item_coords.y,
block_item_coords_k),
block_end_item_k), // block_end_item_k
loader_b(
d_b, // d_matrix
dim_n, // matrix_values_l
(TransformB == matrix_transform_t::NonTranspose) ? 1 : dim_n, // matrix_values_stride_k
(TransformB == matrix_transform_t::NonTranspose) ? dim_k : 1, // matrix_values_stride_l
make_int2( // block_begin_item_coords
grid_raster.block_item_coords.x,
block_item_coords_k),
block_end_item_k), // block_end_item_k
accumulator(scratch->accum_scratch)
{}
//-------------------------------------------------------------------------
// Prefetching utility methods
//-------------------------------------------------------------------------
/**
* Request the calling thread's slices of the shared tiles at depth \p tile_offset_k
*/
inline __device__ void request_local_prefetch(
lds_vector_a_t (&slice_a)[ThreadLdsVectorsA], ///< Slice from A
lds_vector_b_t (&slice_b)[ThreadLdsVectorsB], ///< Slice from B
int tile_offset_k)
{
// Load B strip
for (int i = 0; i < ThreadLdsVectorsB; ++i)
{
slice_b[i].load(
&scratch->pages[page_idx].block_b[tile_offset_k][thread_strip_offset_b + (i * WarpThreadsX * LdsVectorDpVectorsB)]);
}
// Load A strip
for (int i = 0; i < ThreadLdsVectorsA; ++i)
{
slice_a[i].load(
&scratch->pages[page_idx].block_a[tile_offset_k][thread_strip_offset_a + (i * WarpThreadsY * LdsVectorDpVectorsA)]);
}
}
//-------------------------------------------------------------------------
// Epilogue
//-------------------------------------------------------------------------
/**
* Performs the GEMM epilogue:
* - Applies the scalar multipliers and addends to the accumulators
* - Write the result to the output matrix
*/
inline __device__ void epilogue()
{
// Wait for predecessor thread block(s) to produce block-wide tile of
// exclsuive partial-sums
k_split.wait();
// Configure epilogue as to whether the thread block is a secondary
// accumulator in an inter-block k-splitting scheme
if (k_split.is_secondary_accumulator())
epilogue_op.set_secondary_accumulator();
// Whether the addend from C needs loading
bool must_init_addend = epilogue_op.must_init_addend();
#pragma unroll
for (int x = 0; x < ThreadItemsX; ++x)
{
#pragma unroll
for (int y = 0; y < ThreadItemsY; y += LdsVectorDpVectorsA)
{
int thread_strip_b = x / LdsVectorDpVectorsB;
int thread_strip_a = y / LdsVectorDpVectorsA;
int thread_item_coords_tile_x = thread_strip_offset_b + (thread_strip_b * WarpThreadsX * LdsVectorDpVectorsB) + (x % LdsVectorDpVectorsB);
int thread_item_coords_tile_y = thread_strip_offset_a + (thread_strip_a * WarpThreadsY * LdsVectorDpVectorsA) + (y % LdsVectorDpVectorsA);
int c_idx = (grid_raster.block_item_coords.x + thread_item_coords_tile_x) * dim_m +
grid_raster.block_item_coords.y + thread_item_coords_tile_y;
accum_t *my_c = d_c + c_idx;
#pragma unroll
for (int i = 0; i < LdsVectorDpVectorsA; ++i)
{
accum_t c_slice = accum_t(0);
accum_t *c_ptr = my_c + i;
if ((grid_raster.block_item_coords.x + thread_item_coords_tile_x) < dim_n &&
(grid_raster.block_item_coords.y + thread_item_coords_tile_y + i) < dim_m)
{
if (must_init_addend)
{
ldg_cg(c_slice, c_ptr);
}
c_slice = epilogue_op(accumulator.get(x, y + i), c_slice, c_idx + i);
stg_cg(c_ptr, c_slice);
}
}
}
}
// Signal k-split successor thread_block that we have produced our block-wide
// tile of inclusive partial-sums
k_split.signal();
}
//-------------------------------------------------------------------------
// Tile consumption
//-------------------------------------------------------------------------
/**
* Consume a tile of A and B each
*/
template <bool DoGlobalPrefetch>
inline __device__
void consume_tile()
{
// Unroll BlockDpVectorsK iterations of outer-product accumulations
#pragma unroll
for (int tile_offset_k = 0; tile_offset_k < BlockDpVectorsK; tile_offset_k += 1)
{
// Last strip commits global prefetch for next tile
if ((tile_offset_k == BlockDpVectorsK - 1) && DoGlobalPrefetch)
{
// If not using two pages of scratch tiles, protect the above prefetch loads from the committing writes below
if (!UseDoubleScratchTiles)
__syncthreads();
// If using two pages of scratch tiles, switch to next page before writing
if (UseDoubleScratchTiles)
{
page_idx = (page_idx ? 0 : 1);
}
// Commit global prefetch data to scratch page
loader_a.commit(scratch->pages[page_idx].block_a);
loader_b.commit(scratch->pages[page_idx].block_b);
__syncthreads();
}
// Request local prefetch for next strip
request_local_prefetch(
local_slices_a[(tile_offset_k + 1) % 2],
local_slices_b[(tile_offset_k + 1) % 2],
(tile_offset_k + 1) % BlockDpVectorsK);
// Request global prefetch for next tile on first strip
if ((tile_offset_k == 0) && DoGlobalPrefetch)
{
loader_b.request();
loader_b.next();
loader_a.request();
loader_a.next();
}
// Cast strip-mined loads to contiguous array of dp_vector_t
typedef dp_vector_t thread_tile_a_t[ThreadLdsVectorsA * LdsVectorDpVectorsA];
typedef dp_vector_t thread_tile_b_t[ThreadLdsVectorsB * LdsVectorDpVectorsB];
thread_tile_a_t &thread_tile_a = reinterpret_cast<thread_tile_a_t&>(local_slices_a[(tile_offset_k) % 2]);
thread_tile_b_t &thread_tile_b = reinterpret_cast<thread_tile_b_t&>(local_slices_b[(tile_offset_k) % 2]);
// Accumulate this dp-stripe product
accumulator.multiply_accumulate(thread_tile_a, thread_tile_b);
}
}
//-------------------------------------------------------------------------
// GEMM API
//-------------------------------------------------------------------------
/**
* Compute GEMM
*/
inline __device__
void run()
{
// Quit if the thread block is fully out-of-bounds
if (grid_raster.is_block_oob(dim_m, dim_n))
{
asm volatile("exit;");
}
// Request global prefetch of first tile
loader_a.request();
loader_a.next();
loader_b.request();
loader_b.next();
// Commit global prefetch of first tile to shared memory
loader_a.commit(scratch->pages[page_idx].block_a);
loader_b.commit(scratch->pages[page_idx].block_b);
// Advance to next A,B tiles in K-axis
block_item_coords_k += BlockItemsK;
// Synchronize shared tiles and prepared accumulator
__syncthreads();
// Initialize thread's slice of accumulators
accumulator.init();
// Request first iteration of local prefetch strips
request_local_prefetch(
local_slices_a[0],
local_slices_b[0],
0);
//
// Main loop
//
// Consume tiles in A and B along the K-axis (all but last tile)
#pragma unroll 1
while (block_item_coords_k < block_end_item_k)
{
consume_tile<true>();
// Advance to next A,B tiles in K-axis
block_item_coords_k += BlockItemsK;
}
// Consume last tile
consume_tile<false>();
//
// Eplilogue
//
epilogue();
}
};
} // namespace gemm
} // namespace cutlass

View File

@ -1,759 +0,0 @@
/******************************************************************************
* Copyright (c) 2011-2017, NVIDIA CORPORATION. All rights reserved.
*
* Redistribution and use in source and binary forms, with or without
* modification, are not permitted.
*
* THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
* ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
* WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
* DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE FOR ANY
* DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
* (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
* LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
* ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
* (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
* SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
*
******************************************************************************/
/**
* \file
* A block-wide task abstraction for computing device-wide GEMM
*/
#pragma once
// Compiler guard conditional to avoid compilation errors on versions of CUDA that
// do not support the WMMA API.
#if defined (WMMA)
#include <stdint.h>
#include "../util/util.h"
#include "grid_raster.h"
#include "block_loader.h"
#include "block_loader_wmma.h"
#include "wmma_accumulator.h"
namespace cutlass {
namespace gemm {
/******************************************************************************
* block_task_wmma_policy
******************************************************************************/
/**
* \brief Parameterizable tuning policy for block-wide WMMA GEMM tasks
*
* Once parameterized, \p block_task_policy provides the member constant
* \p BlockThreads indicating to the required thread block size
*/
template <
int _BlockItemsY, ///< Height in rows of a block-wide tile in matrix C
int _BlockItemsX, ///< Width in columns of a block-wide tile in matrix C
int _BlockItemsK, ///< Extent of block-wide A|B tiles in value_t along the K-axis
int _WarpItemsY, ///< Height in rows of a Warp tile's accumulators
int _WarpItemsX, ///< Width in columns of a Warp tile's accumulators
int _WmmaItemsY, ///< Height in rows of a discrete WMMA block's accumulators
int _WmmaItemsX, ///< Width in columns of a discrete WMMA block's accumulators
int _WmmaItemsK, ///< Depth of each discrete WMMA block
bool _UseDoubleScratchTiles, ///< Whether to halve synchronization overhead at the expense of doubled shared memory and addressing overhead
grid_raster_strategy::kind_t _RasterStrategy> ///< Strategy for enumerating \p block_task within an input matrix
struct block_task_wmma_policy
{
/// Strategy for enumerating \p block_task within an input matrix
static const grid_raster_strategy::kind_t RasterStrategy = _RasterStrategy;
enum
{
/// Height in rows of a block-wide tile in matrix C
BlockItemsY = _BlockItemsY,
/// Width in columns of a block-wide tile in matrix C
BlockItemsX = _BlockItemsX,
/// Extent of block-wide A|B tiles in value_t along the K-axis
BlockItemsK = _BlockItemsK,
/// Height in rows of a Warp tile's accumulators
WarpItemsX = _WarpItemsX,
/// Width in columns of a Warp tile's accumulators
WarpItemsY = _WarpItemsY,
/// Width in columns of a discrete WMMA block's accumulators
WmmaItemsX = _WmmaItemsX,
/// Height in rows of a discrete WMMA block's accumulators
WmmaItemsY = _WmmaItemsY,
/// Depth of each discrete WMMA block
WmmaItemsK = _WmmaItemsK,
/// Whether to halve synchronization overhead at the expense of doubled shared memory and addressing overhead
UseDoubleScratchTiles = _UseDoubleScratchTiles,
//
// Derived quantities
//
/// Machine warp size
WarpThreads = 32,
/// Number of WMMA operations in the height dimension
WmmaBlocksY = divide_assert<WarpItemsY, WmmaItemsY>::value,
/// Number of WMMA operations in the height dimension
WmmaBlocksX = divide_assert<WarpItemsX, WmmaItemsX>::value,
/// Number of warps in each thread block
BlockWarps = divide_assert<BlockItemsY * BlockItemsX, WarpItemsX * WarpItemsY>::value,
/// Number of threads in each thread block (blockDim.x)
BlockThreads = BlockWarps * WarpThreads,
};
};
/******************************************************************************
* block_task_wmma
******************************************************************************/
/**
* \brief A block-wide task abstraction for computing device-wide GEMM
*
* Each thread_block is assigned a unique tile of output matrix C to compute by
* consuming the corresponding stripes of the input matrices A and B.
*/
template <
typename block_task_policy_t, ///< Parameterization of block_task_policy
typename value_t, ///< Multiplicand value type (matrices A and B)
typename accum_t, ///< Accumulator value type (matrix C and scalars)
matrix_transform_t::kind_t TransformA, ///< View transform enumerant for matrix A
int LdgAlignA, ///< Alignment (in bytes) for A operand
matrix_transform_t::kind_t TransformB, ///< View transform enumerant for matrix B
int LdgAlignB, ///< Alignment (in bytes) for B operand
typename epilogue_op_t, ///< Epilogue operation to update matrix C
int LdgAlignC, ///< Alignment (in bytes) for C operand
bool AllowRaggedTiles ///< Whether the input matrix's dimensions need not be an even-multiple of the block-wide tile dimensions
>
struct block_task_wmma
{
//-------------------------------------------------------------------------
// Constants and types
//-------------------------------------------------------------------------
enum
{
/// Number of threads in each thread block (blockDim.x)
BlockThreads = block_task_policy_t::BlockThreads,
/// Extent of block-wide C-tile in accum_t (and A-tiles in value_t) along M-axis (height)
BlockItemsY = block_task_policy_t::BlockItemsY,
/// Extent of block-wide C-tile in accum_t (and B-tiles in value_t) along N-axis (width)
BlockItemsX = block_task_policy_t::BlockItemsX,
/// Extent of block-wide A|B tiles in value_t along the K-axis
BlockItemsK = block_task_policy_t::BlockItemsK,
/// Extent of warp C-tile in accum_t (and A-tiles in value_t) along M-axis (height)
WarpItemsY = block_task_policy_t::WarpItemsY,
/// Extent of warp C-tile in accum_t (and B-tiles in value_t) along N-axis (width)
WarpItemsX = block_task_policy_t::WarpItemsX,
/// Extent of warp C-tile in accum_t (and A-tiles in value_t) along M-axis (height)
WmmaItemsY = block_task_policy_t::WmmaItemsY,
/// Extent of warp C-tile in accum_t (and B-tiles in value_t) along N-axis (width)
WmmaItemsX = block_task_policy_t::WmmaItemsX,
/// Extent of warp-wide A|B-tiles in value_t along K-axis
WmmaItemsK = block_task_policy_t::WmmaItemsK,
/// Whether to halve synchronization overhead at the expense of doubled shared memory and addressing overhead
UseDoubleScratchTiles = block_task_policy_t::UseDoubleScratchTiles,
/// Number of threads in warp
WarpThreads = block_task_policy_t::WarpThreads,
/// Number of warps participating
BlockWarps = block_task_policy_t::BlockWarps,
/// Extent of block in warps along M-axis
BlockWarpsY = divide_assert<BlockItemsY, WarpItemsY>::value,
/// Extent of block in warps along N-axis
BlockWarpsX = divide_assert<BlockItemsX, WarpItemsX>::value,
/// Number of MMA unrolls
WmmaUnrollCount = divide_assert<BlockItemsK, WmmaItemsK>::value,
/// True if the A matrix layout is column major (K is the strided dimension)
IsLayoutCongruousA = (TransformA == matrix_transform_t::NonTranspose),
/// True if the B matrix layout is row mayor (K is the strided dimension)
IsLayoutCongruousB = (TransformB == matrix_transform_t::Transpose),
};
/// WMMA may support unique types for A and B, so plan ahead for this
typedef value_t value_a_t;
/// WMMA may support unique types for A and B, so plan ahead for this
typedef value_t value_b_t;
/// WMMA accumulator type
typedef wmma_accumulator<
WarpItemsY,
WarpItemsX,
WmmaItemsY,
WmmaItemsX,
WmmaItemsK,
value_a_t,
value_b_t,
accum_t,
TransformA,
TransformB>
accumulator_t;
/// Thread block rasterization helper type
typedef grid_raster<
BlockItemsY,
BlockItemsX,
TransformA,
TransformB,
block_task_policy_t::RasterStrategy>
grid_raster_t;
/// Tile loader type for matrix A
typedef block_loader_wmma<
IsLayoutCongruousA,
BlockThreads,
(IsLayoutCongruousA ? BlockItemsY : BlockItemsK),
(IsLayoutCongruousA ? BlockItemsK : BlockItemsY),
value_a_t,
LdgAlignA,
AllowRaggedTiles>
block_loader_a_t;
/// Tile loader type for matrix A
typedef block_loader_wmma<
IsLayoutCongruousB,
BlockThreads,
(IsLayoutCongruousB ? BlockItemsX : BlockItemsK),
(IsLayoutCongruousB ? BlockItemsK : BlockItemsX),
value_b_t,
LdgAlignB,
AllowRaggedTiles>
block_loader_b_t;
/// Type alias for matrix A fragment type
typedef typename accumulator_t::fragment_a_t fragment_a_t;
/// Type alias for matrix B fragment type
typedef typename accumulator_t::fragment_b_t fragment_b_t;
enum
{
/// Number of fragments from A matrix
WmmaBlocksY = accumulator_t::WmmaBlocksY,
/// Number of fragments from B matrix
WmmaBlocksX = accumulator_t::WmmaBlocksX,
/// Number of value_t to pad the outer dimension of the shared A-tile
PadItemsA = 16,
/// Number of value_t to pad the outer dimension of the shared B-tile
PadItemsB = 16,
/// Leading dimension of A matrix tile
LdmSmemA = (IsLayoutCongruousA ? BlockItemsY: BlockItemsK) + PadItemsA,
/// Leading dimension of A matrix tile
StridedSmemA = (IsLayoutCongruousA ? BlockItemsK : BlockItemsY ),
/// Leading dimension of B matrix tile
LdmSmemB = (IsLayoutCongruousB? BlockItemsX : BlockItemsK) + PadItemsB,
StridedSmemB = (IsLayoutCongruousB ? BlockItemsK : BlockItemsX),
};
/// Shared memory layout for a prefetch page
struct page_storage_t
{
/// Tile of A
value_a_t __align__(16) block_a[StridedSmemA][LdmSmemA];
/// Tile of B
value_b_t __align__(16) block_b[StridedSmemB][LdmSmemB];
};
/// Shared memory layout for scratch storage
struct scratch_storage_t
{
union
{
/// Prefetch pages
uninitialized<page_storage_t> pages[UseDoubleScratchTiles ? 2 : 1];
/// Scratch storage for warps
accum_t epilogue[BlockWarps][WmmaItemsX * WmmaItemsY];
};
};
//-------------------------------------------------------------------------
// Assert assumptions
//-------------------------------------------------------------------------
// Ensure we have at least two unrolled innermost loop iterations (one to prefetch
// the next global tile and then one to prefetch the first strip of it from shared)
static_assert ((BlockItemsK >= 2), "BlockItemsK must be >= 2.");
//-------------------------------------------------------------------------
// Members
//-------------------------------------------------------------------------
/// Scratch storage reference
scratch_storage_t *scratch;
/// Which page of scratch tiles we're currently reading from
int page_idx;
/// Pointer to matrix C
accum_t *d_c;
/// Epilogue operation applied to update matrix C
epilogue_op_t epilogue_op;
/// Matrix height in rows of trans_op(A) and C
int dim_m;
/// Matrix width in columns of trans_op(B) and C
int dim_n;
/// Control for inter-block k-splitting
k_split_control k_split;
/// Thread block's base value_t coordinates (m, n) in matrix C
grid_raster_t grid_raster;
/// Thread block's current coordinate (k) within A|B matrices
int block_item_coords_k;
/// Thread block's ending coordinate (k) within A|B matrices (one-past)
int block_end_item_k;
/// Warp's coordinates (x, y) in thread block
int2 block_warp_item_coords;
/// A tile loader
block_loader_a_t loader_a;
/// B tile loader
block_loader_b_t loader_b;
/// Thread's active-k/prefetch-k slices from shared A tile
fragment_a_t local_slices_a[2][WmmaBlocksY];
/// Thread's active-k/prefetch-k slices from shared B tile
fragment_b_t local_slices_b[2][WmmaBlocksX];
/// Accumulator tile
accumulator_t accumulator;
//-------------------------------------------------------------------------
// Coordinate system helpers
//-------------------------------------------------------------------------
/// Compute the warp's item-coordinates (x, y) in thread block
inline __device__
int2 warp_item_coords()
{
int warp_id = threadIdx.x / WarpThreads;
return make_int2(
(warp_id / BlockWarpsY) * WarpItemsX,
(warp_id % BlockWarpsY) * WarpItemsY);
}
/// Compute the thread block's base item-coordinates in matrix A
inline __device__
int2 a_block_item_coords()
{
if (TransformA == matrix_transform_t::NonTranspose)
{
return make_int2(grid_raster.block_item_coords.y, block_item_coords_k);
}
else
{
return make_int2(block_item_coords_k, grid_raster.block_item_coords.y);
}
}
/// Compute the thread block's base item-coordinates in matrix B
inline __device__
int2 b_block_item_coords()
{
if (TransformB == matrix_transform_t::Transpose)
{
return make_int2(grid_raster.block_item_coords.x, block_item_coords_k);
}
else
{
return make_int2(block_item_coords_k, grid_raster.block_item_coords.x);
}
}
//-------------------------------------------------------------------------
// Constructor API
//-------------------------------------------------------------------------
/// Constructor
inline __device__
block_task_wmma(
scratch_storage_t *scratch,
value_t *d_a,
value_t *d_b,
accum_t *d_c,
epilogue_op_t epilogue_op,
int dim_m,
int dim_n,
int dim_k,
k_split_control k_split)
:
scratch(scratch),
page_idx(0),
d_c(d_c),
epilogue_op(epilogue_op),
dim_m(dim_m),
dim_n(dim_n),
k_split(k_split),
block_item_coords_k(k_split.block_begin_item_k()),
block_end_item_k(k_split.block_end_item_k(dim_k)),
block_warp_item_coords(warp_item_coords()),
loader_a(
reinterpret_cast<value_a_t const *>(d_a),
(IsLayoutCongruousA ? dim_m : block_end_item_k),
(IsLayoutCongruousA ? 0 : block_item_coords_k),
(IsLayoutCongruousA ? block_end_item_k : dim_m),
(IsLayoutCongruousA ? dim_m : dim_k),
(IsLayoutCongruousA ? block_item_coords_k : 0),
a_block_item_coords()),
loader_b(
reinterpret_cast<value_b_t const *>(d_b),
(IsLayoutCongruousB ? dim_n : block_end_item_k),
(IsLayoutCongruousB ? 0 : block_item_coords_k),
(IsLayoutCongruousB ? block_end_item_k : dim_n),
(IsLayoutCongruousB ? dim_n : dim_k),
(IsLayoutCongruousB ? block_item_coords_k : 0),
b_block_item_coords())
{}
//-------------------------------------------------------------------------
// Prefetching utility methods
//-------------------------------------------------------------------------
/**
* Request the calling thread's slices of the shared tiles at depth \p tile_offset_k
*/
inline __device__ void request_local_prefetch(
fragment_a_t local_slices_a[WmmaBlocksY], ///< Slice from A
fragment_b_t local_slices_b[WmmaBlocksX], ///< Slice from B
int tile_offset_k)
{
value_b_t const *smem_A_base = &scratch->pages[page_idx].alias().block_a[0][0];
value_b_t const *smem_B_base = &scratch->pages[page_idx].alias().block_b[0][0];
int constexpr kstride_a = (IsLayoutCongruousA ? LdmSmemA : 1);
int constexpr lstride_a = (IsLayoutCongruousA ? 1 : LdmSmemA);
int constexpr kstride_b = (IsLayoutCongruousB ? LdmSmemB : 1);
int constexpr lstride_b = (IsLayoutCongruousB ? 1 : LdmSmemB);
// Load B strip
#pragma unroll
for (int i = 0; i < WmmaBlocksX; ++i)
{
value_b_t const *smem_B_ptr =
&smem_B_base[tile_offset_k * kstride_b + (block_warp_item_coords.x + WmmaItemsX * i) * lstride_b];
nvcuda::wmma::load_matrix_sync(local_slices_b[i], smem_B_ptr, LdmSmemB);
}
// Load A strip
#pragma unroll
for (int i = 0; i < WmmaBlocksY; ++i)
{
value_a_t const *smem_A_ptr =
&smem_A_base[tile_offset_k * kstride_a + (block_warp_item_coords.y + WmmaItemsY * i) * lstride_a];
nvcuda::wmma::load_matrix_sync(local_slices_a[i], smem_A_ptr, LdmSmemA);
}
}
//-------------------------------------------------------------------------
// Epilogue
//-------------------------------------------------------------------------
/**
* Performs the GEMM epilogue:
* - Applies the scalar multipliers and addends to the accumulators
* - Write the result to the output matrix
*/
inline __device__ void epilogue()
{
// Wait for predecessor thread block(s) to produce partial-sums
k_split.wait();
// Configure epilogue as to whether the thread block is a secondary
// accumulator in an inter-block k-splitting scheme
if (k_split.is_secondary_accumulator())
epilogue_op.set_secondary_accumulator();
// Whether or not the addend from C needs loading
bool must_init_addend = epilogue_op.must_init_addend();
int warp_base_x = grid_raster.block_item_coords.x + block_warp_item_coords.x;
int warp_base_y = grid_raster.block_item_coords.y + block_warp_item_coords.y;
int constexpr SmemStride = WmmaItemsY;
int warp_id = threadIdx.x / 32;
// Compute shape of one accumulator read/modify/write operation
int constexpr ItemsY = (WmmaItemsY);
int constexpr ItemsX = (32 / ItemsY);
int constexpr IterationsX = WmmaItemsX / ItemsX;
// Compute a rasterization of warp lanes across the WMMA tile.
int lane_id = (threadIdx.x % 32);
int lane_read_x = (lane_id / ItemsY);
int lane_read_y = (lane_id % ItemsY);
accum_t *smem_scratch = scratch->epilogue[warp_id];
accum_t const *smem_read_ptr = smem_scratch + lane_read_y + lane_read_x * SmemStride;
#pragma unroll
for (int xb = 0; xb < WmmaBlocksX; ++xb)
{
#pragma unroll
for (int yb = 0; yb < WmmaBlocksY; ++yb)
{
// Store accumulator tile to SMEM
nvcuda::wmma::store_matrix_sync(
smem_scratch,
accumulator.accumulators[xb][yb],
SmemStride,
matrix_layout<matrix_transform_t::NonTranspose>::kind);
// Synchronize threads within the warp
__syncthreads();
// Compute lane coordinates so that each thread efficiently accesses SMEM.
int c_x = (warp_base_x + (xb) * WmmaItemsX + lane_read_x);
int c_y = (warp_base_y + (yb) * WmmaItemsY + lane_read_y);
// Compute guard predicate by comparing against problem dimensions.
bool pred = c_y < dim_m;
// Compute output pointer from lane coordinates
int c_index = c_x * dim_m + c_y;
accum_t *c_ptr = reinterpret_cast<accum_t *>(d_c) + c_x * dim_m + c_y;
// Iterate over columns of output tile. Load from SMEM, compute epilogue operation,
// and stream output to global memory
#pragma unroll
for (int item_x = 0; item_x < IterationsX; ++item_x)
{
accum_t accum = smem_read_ptr[item_x * ItemsX * SmemStride];
accum_t c_element = 0;
// Filter against problem dimensions as the warp iterates across the columns of
// output.
pred = (pred && ((c_x + item_x * ItemsX) < dim_n));
if (must_init_addend && pred)
{
// NB: inline PTX to utilize strong operations for inter-block synchronization.
// The following is equivalent to:
//
// c_element = c_ptr[0];
asm volatile ("ld.global.cg.f32 %0, [%1];\n" : "=f"(c_element) : "l"(c_ptr));
}
c_element = epilogue_op(accum, c_element, c_index);
if (pred)
{
// NB: inline PTX to utilize strong operations for inter-block synchronization.
// The following is equivalent to:
//
// c_ptr[0] = c_element;
asm volatile ("st.global.cg.f32 [%0], %1;\n" : : "l"(c_ptr), "f"(c_element));
}
// Increment output pointer
c_ptr += dim_m * ItemsX;
c_index += dim_m * ItemsX;
}
__syncthreads();
}
}
// Signal k-split successor thread_block
k_split.signal();
}
//-------------------------------------------------------------------------
// Tile consumption
//-------------------------------------------------------------------------
/**
* Consume a tile of A and B each
*/
template <bool DoGlobalPrefetch>
inline __device__
void consume_tile()
{
// Request global prefetch for next tile on first strip
if (DoGlobalPrefetch)
{
loader_b.request();
loader_b.next();
loader_a.request();
loader_a.next();
}
// Unroll BlockDpVectorsK iterations of outer-product accumulations
#pragma unroll
for (int iteration = 0; iteration < WmmaUnrollCount; ++iteration)
{
int tile_offset_k = iteration * WmmaItemsK;
// Active load-from-shared index
int active_lds_idx = __NV_STD_MIN(WmmaUnrollCount - 1, (iteration) % 2);
// Next load-from-shared index
int next_lds_idx = __NV_STD_MIN(WmmaUnrollCount - 1, (iteration + 1) % 2);
// The last unrolled iteration commits the global fetches
if ((iteration == WmmaUnrollCount - 1) && DoGlobalPrefetch)
{
// If not using two pages of scratch tiles, protect the above prefetch loads from
// the committing writes below
if (!UseDoubleScratchTiles)
{
__syncthreads();
}
else
{
page_idx = (page_idx ? 0 : 1);
}
// Commit global prefetch data to scratch page
loader_a.template commit<LdmSmemA>(&scratch->pages[page_idx].alias().block_a[0][0]);
loader_b.template commit<LdmSmemB>(&scratch->pages[page_idx].alias().block_b[0][0]);
__syncthreads();
}
// Accumulate this dp-stripe product
accumulator.multiply_accumulate(
local_slices_a[active_lds_idx],
local_slices_b[active_lds_idx]);
// Request local prefetch for next strip
request_local_prefetch(
local_slices_a[next_lds_idx],
local_slices_b[next_lds_idx],
(tile_offset_k + WmmaItemsK) % BlockItemsK);
}
}
//-------------------------------------------------------------------------
// GEMM API
//-------------------------------------------------------------------------
/**
* Compute GEMM
*/
inline __device__
void run()
{
// Quit if the thread block is fully out-of-bounds
if (grid_raster.is_block_oob(dim_m, dim_n))
{
asm volatile("exit;");
}
// Request global prefetch of first tile
loader_a.request();
loader_a.next();
loader_b.request();
loader_b.next();
// Commit global prefetch of first tile to shared memory
loader_a.template commit<LdmSmemA>(&scratch->pages[page_idx].alias().block_a[0][0]);
loader_b.template commit<LdmSmemB>(&scratch->pages[page_idx].alias().block_b[0][0]);
// Advance to next A,B tiles in K-axis
block_item_coords_k += BlockItemsK;
// Synchronize shared tiles and prepared accumulator
__syncthreads();
// Initialize thread's slice of accumulators
accumulator.init();
// Request first iteration of local prefetch strips
request_local_prefetch(
local_slices_a[0],
local_slices_b[0],
0);
//
// Main loop
//
// Consume tiles in A and B along the K-axis (all but last tile)
#pragma unroll 1
while (block_item_coords_k < block_end_item_k)
{
consume_tile<true>();
// Advance to next A,B tiles in K-axis
block_item_coords_k += BlockItemsK;
}
consume_tile<false>();
//
// Eplilogue
//
// prevent overwriting SMEM until all warps have finished loading data
__syncthreads();
// store accumulator tile to global memory
epilogue();
}
};
} // namespace gemm
} // namespace cutlass
#endif

View File

@ -1,534 +0,0 @@
/******************************************************************************
* Copyright (c) 2011-2017, NVIDIA CORPORATION. All rights reserved.
*
* Redistribution and use in source and binary forms, with or without
* modification, are not permitted.
*
* THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
* ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
* WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
* DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE FOR ANY
* DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
* (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
* LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
* ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
* (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
* SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
*
******************************************************************************/
#pragma once
/**
* \file
* GEMM kernel entrypoint and dispatch stub
*/
#include <stdint.h>
#include "../util/util.h"
#include "block_task.h"
#include "block_task_wmma.h"
#include "grid_raster.h"
#include "dispatch_policies.h"
#include "k_split_control.h"
namespace cutlass {
namespace gemm {
/******************************************************************************
* param_pack
******************************************************************************/
/**
* Parameter-pack structure
*
* Kernel launch latency is reduced when kernel arguments are wrapped into
* a single parameter
*/
template <
typename value_t,
typename accum_t,
typename epilogue_op_t>
struct param_pack
{
int m; ///< Height in rows of op(A) and C
int n; ///< Width in columns of op(B) and C
int k; ///< Width in columns of op(A) and height in rows of op(B)
k_split_control k_split; ///< Abstraction for controlling inter-block k-splitting
value_t *d_a; ///< Pointer to matrix A array values
value_t *d_b; ///< Pointer to matrix B array values
accum_t *d_c; ///< Pointer to matrix C array values
epilogue_op_t epilogue_op;
param_pack(
int m, ///< Height in rows of op(A) and C
int n, ///< Width in columns of op(B) and C
int k, ///< Width in columns of op(A) and height in rows of op(B)
k_split_control k_split, ///< Abstraction for controlling inter-block k-splitting
epilogue_op_t op, ///< Epilogue operation to update matrix C
value_t *d_a, ///< Pointer to matrix A array values
value_t *d_b, ///< Pointer to matrix B array values
accum_t *d_c) ///< Pointer to matrix C array values
:
m(m),
n(n),
k(k),
k_split(k_split),
epilogue_op(op),
d_a(d_a),
d_b(d_b),
d_c(d_c)
{}
};
/******************************************************************************
* Conditionally select the appropriate GEMM threadblock task
******************************************************************************/
/// Conditional selection for block task
template <
math_operation_class_t math_op, ///<
typename block_task_policy_t, ///< Parameterization of block_task_policy
typename value_t, ///< Multiplicand value type (matrices A and B)
typename accum_t, ///< Accumulator value type (matrix C and scalars)
matrix_transform_t::kind_t TransformA, ///< View transform enumerant for matrix A
int LdgAlignA, ///< Alignment (in bytes) for A operand
matrix_transform_t::kind_t TransformB, ///< View transform enumerant for matrix B
int LdgAlignB, ///< Alignment (in bytes) for B operand
typename epilogue_op_t, ///< Epilogue operation applied to GEMM
int LdgAlignC, ///< Alignment (in bytes) for C operand
bool AllowRaggedTiles ///< Whether GEMM supports matrix sizes other than multiple of BlockItems{XY}
>
struct gemm_block_task;
/// Scalar math operations
template <
typename block_task_policy_t, ///< Parameterization of block_task_policy
typename value_t, ///< Multiplicand value type (matrices A and B)
typename accum_t, ///< Accumulator value type (matrix C and scalars)
matrix_transform_t::kind_t TransformA, ///< View transform enumerant for matrix A
int LdgAlignA, ///< Alignment (in bytes) for A operand
matrix_transform_t::kind_t TransformB, ///< View transform enumerant for matrix B
int LdgAlignB, ///< Alignment (in bytes) for B operand
typename epilogue_op_t, ///< Epilogue operation applied to GEMM
int LdgAlignC, ///< Alignment (in bytes) for C operand
bool AllowRaggedTiles ///< Whether GEMM supports matrix sizes other than multiple of BlockItems{XY}
>
struct gemm_block_task<
math_operation_class_t::scalar,
block_task_policy_t,
value_t,
accum_t,
TransformA,
LdgAlignA,
TransformB,
LdgAlignB,
epilogue_op_t,
LdgAlignC,
AllowRaggedTiles
>
{
// Parameterize task type
typedef block_task<
block_task_policy_t,
value_t,
accum_t,
TransformA,
LdgAlignA,
TransformB,
LdgAlignB,
epilogue_op_t,
LdgAlignC,
AllowRaggedTiles> type;
};
/// Matrix math operations
template <
typename block_task_policy_t, ///< Parameterization of block_task_policy
typename value_t, ///< Multiplicand value type (matrices A and B)
typename accum_t, ///< Accumulator value type (matrix C and scalars)
matrix_transform_t::kind_t TransformA, ///< View transform enumerant for matrix A
int LdgAlignA, ///< Alignment (in bytes) for A operand
matrix_transform_t::kind_t TransformB, ///< View transform enumerant for matrix B
int LdgAlignB, ///< Alignment (in bytes) for B operand
typename epilogue_op_t, ///< Epilogue operation applied to GEMM
int LdgAlignC, ///< Alignment (in bytes) for C operand
bool AllowRaggedTiles ///< Whether GEMM supports matrix sizes other than multiple of BlockItems{XY}
>
struct gemm_block_task<
math_operation_class_t::matrix,
block_task_policy_t,
value_t,
accum_t,
TransformA,
LdgAlignA,
TransformB,
LdgAlignB,
epilogue_op_t,
LdgAlignC,
AllowRaggedTiles>
{
#if defined(WMMA) // conditional compilation with WMMA headers
// Parameterize task type
typedef block_task_wmma<
block_task_policy_t,
value_t,
accum_t,
TransformA,
LdgAlignA,
TransformB,
LdgAlignB,
epilogue_op_t,
LdgAlignC,
AllowRaggedTiles> type;
#endif
};
/******************************************************************************
* GEMM kernel entrypoint
******************************************************************************/
/**
* GEMM kernel
*
* NB: Not sure why NVVM is doing stuff with "__launch_bounds__" instead of just
* passing it along to PTXAS, but it is currently resulting in less optimal codegen
*/
template <
math_operation_class_t math_op, ///< Indicates which class of math operation to select
typename block_task_policy_t, ///< Parameterization of block_task_policy
matrix_transform_t::kind_t TransformA, ///< Transformation op for matrix A
int LdgAlignA, ///< Alignment of A matrix elements in bytes
matrix_transform_t::kind_t TransformB, ///< Transformation op for matrix B
int LdgAlignB, ///< Alignment of B matrix elements in bytes
typename value_t, ///< Multiplicand value type (matrices A and B)
typename accum_t, ///< Accumulator value type (matrix C and scalars)
typename epilogue_op_t, ///< Epilogue operation applied to update matrix C
int LdgAlignC, ///< Alignment of C elements in bytes
bool AllowRaggedTiles> ///< Boolean to indicate whether AllowRaggedTiles handling is enabled
__global__ void kernel(param_pack<value_t, accum_t, epilogue_op_t> pack)
{
// Parameterize task type
typedef typename gemm_block_task<
math_op,
block_task_policy_t,
value_t,
accum_t,
TransformA,
LdgAlignA,
TransformB,
LdgAlignB,
epilogue_op_t,
LdgAlignC,
AllowRaggedTiles>::type block_task_t;
// Declare statically-allocated shared storage
__shared__ typename block_task_t::scratch_storage_t smem;
// Construct and run the task
block_task_t(
&smem,
pack.d_a,
pack.d_b,
pack.d_c,
pack.epilogue_op,
pack.m,
pack.n,
pack.k,
pack.k_split).run();
}
/******************************************************************************
* Launch configuration description returned to the caller
******************************************************************************/
/// Return details about the launch configuration to the caller
struct launch_configuration
{
//
// Data members
//
/// cudaError_t resulting from grid launch
cudaError_t result;
/// Extent of a thread block's partition along the GEMM K-axis
int split_k;
/// Kernel grid extents in thread blocks
dim3 grid;
/// Thread block extents in threads
dim3 block;
//
// Methods
//
/// Constructor
launch_configuration():
result(cudaSuccess),
split_k(0),
grid(0, 0, 0),
block(0, 0, 0) {
}
/// Conversion from cudaError_t
launch_configuration(cudaError_t result):
result(result),
split_k(1),
grid(0, 0, 0),
block(0, 0, 0) {
}
/// Launch configuration for Cutlass kernels
launch_configuration(
cudaError_t result,
int split_k,
dim3 grid,
dim3 block
):
result(result),
split_k(split_k),
grid(grid),
block(block) {
}
};
/******************************************************************************
* Dispatch stub
******************************************************************************/
/**
* GEMM dispatch stub
*
* This function also serves as the autotuning entrypoint to evaluate different
* tuning parameterizations of kernel.
*/
template <
math_operation_class_t math_op, ///< Indicates which class of math operation to select
typename block_task_policy_t, ///< Parameterization of block_task_policy
matrix_transform_t::kind_t TransformA, ///< Transformation op for matrix A
int LdgAlignA, ///< Alignment of A matrix elements in bytes
matrix_transform_t::kind_t TransformB, ///< Transformation op for matrix B
int LdgAlignB, ///< Alignment of B matrix elements in bytes
typename value_t, ///< Multiplicand value type (matrices A and B)
typename accum_t, ///< Accumulator value type (matrix C and scalars)
typename epilogue_op_t, ///< Epilogue operation
int LdgAlignC, ///< Alignment of C matrix elements in bytes
bool AllowRaggedTiles, ///< Boolean to indicate whether AllowRaggedTiles handling is enabled
typename kernel_ptr_t> ///< GEMM kernel function pointer type
launch_configuration dispatch(
kernel_ptr_t kernel_ptr, ///< GEMM kernel function pointer
int m, ///< Height in rows of op(A) and C
int n, ///< Width in columns of op(B) and C
int k, ///< Width in columns of op(A) and height in rows of op(B)
epilogue_op_t epilogue_op, ///< Epilogue operation to update matrix C
value_t *d_a, ///< Device pointer to matrix A array values
value_t *d_b, ///< Device pointer to matrix B array values
accum_t *d_c, ///< Device pointer to matrix C array values
cudaStream_t stream = 0, ///< CUDA stream to launch kernels within. Default is stream<sub>0</sub>.
bool debug_synchronous = true) ///< Whether or not to synchronize the stream after every kernel launch
/// to check for errors. Also causes launch configurations to be printed
/// to the console if DEBUG is defined. Default is \p false.
{
// Thread block rasterization type
typedef grid_raster<
block_task_policy_t::BlockItemsY,
block_task_policy_t::BlockItemsX,
TransformA,
TransformB,
block_task_policy_t::RasterStrategy>
grid_raster_t;
launch_configuration config;
// Compute block dims
config.block = dim3(block_task_policy_t::BlockThreads);
// Compute shared memory
int dynamic_smem_bytes = 0;
// Compute occupancy
int max_sm_occupancy;
if (CUDA_PERROR_DEBUG(config.result = cudaOccupancyMaxActiveBlocksPerMultiprocessor(
&max_sm_occupancy,
kernel_ptr,
config.block.x * config.block.y,
dynamic_smem_bytes)))
{
return config;
}
// Compute grid extents
config.grid = grid_raster_t::grid_dims(m, n);
// Get SM count
int sm_count;
if (CUDA_PERROR_DEBUG(config.result = get_sm_count(sm_count)))
return config;
// Get k-split flag storage (TODO: make a pool)
int *d_flags;
if (CUDA_PERROR_DEBUG(config.result = cudaGetSymbolAddress((void**) &d_flags, d_flags_split_k)))
return config;
// Construct k-split coordinator
k_split_control k_split(
d_flags,
sm_count,
max_sm_occupancy,
k,
block_task_policy_t::BlockItemsK,
config.block,
config.grid); // in,out
config.split_k = k_split.split_k;
// Log kernel configuration
if (debug_synchronous)
{
// Compute tiling efficiency
float block_tiling_efficiency = float(block_task_policy_t::BlockItemsY * block_task_policy_t::BlockItemsX) /
float(block_task_policy_t::BlockItemsY + block_task_policy_t::BlockItemsX);
float tiling_efficiency = block_tiling_efficiency;
float wave_efficiency = k_split.get_wave_efficiency(
sm_count, max_sm_occupancy, config.block, config.grid);
CUDA_LOG_DEBUG("Final wave_efficiency %.4f, tiling_efficiency %.4f\n",
wave_efficiency, tiling_efficiency);
CUDA_LOG_DEBUG("Invoking kernel<<<(%d, %d, %d), (%d.y,%d.x), %d, %lld>>>(), %d SM occupancy, %d split_k\n",
config.grid.x, config.grid.y, config.grid.z,
config.block.y, config.block.x,
dynamic_smem_bytes,
(long long) stream,
max_sm_occupancy,
k_split.split_k);
}
// Construct parameter-pack
param_pack<value_t, accum_t, epilogue_op_t> pack(
m,
n,
k,
k_split,
epilogue_op,
d_a,
d_b,
d_c);
// Prepare k-split coordinator
if (CUDA_PERROR_DEBUG(config.result = k_split.prepare(stream, debug_synchronous)))
{
return config;
}
// Invoke kernel
kernel_ptr<<< config.grid, config.block, dynamic_smem_bytes, stream >>>(pack);
// Check for failure to launch
if (CUDA_PERROR_DEBUG(config.result = cudaPeekAtLastError()))
return config;
// Sync the stream if specified to flush runtime errors
if (debug_synchronous && (CUDA_PERROR_DEBUG(config.result = cudaStreamSynchronize(stream))))
return config;
return config;
}
/******************************************************************************
* GEMM
******************************************************************************/
/**
* Computes gemm on device matrices
*/
template <
tiling_strategy::kind_t TilingStrategy, ///< Tile-sizing classification
math_operation_class_t math_op, ///< Indicates which class of math operation to select
matrix_transform_t::kind_t TransformA, ///< Transformation op for matrix A
int LdgAlignA, ///< Alignment (in bytes) of A operand
matrix_transform_t::kind_t TransformB, ///< Transformation op for matrix B
int LdgAlignB, ///< Alignment (in bytes) of B operand
typename value_t, ///< Multiplicand value type (matrices A and B)
typename accum_t, ///< Accumulator value type (matrix C and scalars)
typename epilogue_op_t, ///< Epilogue operation to update matrix C
int LdgAlignC> ///< Alignment (in bytes) of C operand
launch_configuration device_gemm(
int m, ///< Height in rows of op(A) and C
int n, ///< Width in columns of op(B) and C
int k, ///< Width in columns of op(A) and height in rows of op(B)
epilogue_op_t epilogue_op, ///< Epilogue operation to update matrix C
value_t *d_a, ///< Device pointer to matrix A array values
value_t *d_b, ///< Device pointer to matrix B array values
accum_t *d_c, ///< Device pointer to matrix C array values
cudaStream_t stream = 0, ///< CUDA stream to launch kernels within. Default is stream<sub>0</sub>.
bool debug_synchronous = false) ///< Whether or not to synchronize the stream after every kernel launch to
/// check for errors. Also causes launch configurations to be printed to
/// the console if DEBUG is defined. Default is \p false.
{
// Parameterize an task policy type
// (TODO: use a policy dispatch mechanism based upon SM version)
typedef gemm_policy<value_t, accum_t, TransformA, TransformB, TilingStrategy> block_task_policy_t;
// AllowRaggedTiles-tile check
if ((m % block_task_policy_t::BlockItemsY != 0) ||
(n % block_task_policy_t::BlockItemsX != 0) ||
(k % block_task_policy_t::BlockItemsK != 0))
{
// Needs ragged tile-handling
static const bool AllowRaggedTiles = true;
return dispatch<math_op, block_task_policy_t, TransformA, LdgAlignA, TransformB, LdgAlignB, value_t, accum_t, epilogue_op_t, LdgAlignC, AllowRaggedTiles>(
kernel<math_op,block_task_policy_t, TransformA, LdgAlignA, TransformB, LdgAlignB, value_t, accum_t, epilogue_op_t, LdgAlignC, AllowRaggedTiles>,
m,
n,
k,
epilogue_op,
d_a,
d_b,
d_c,
stream,
debug_synchronous);
}
else
{
// Does not need ragged tile-handling
static const bool AllowRaggedTiles = false;
return dispatch<math_op, block_task_policy_t, TransformA, LdgAlignA, TransformB, LdgAlignB, value_t, accum_t, epilogue_op_t, LdgAlignC, AllowRaggedTiles>(
kernel<math_op,block_task_policy_t, TransformA, LdgAlignA, TransformB, LdgAlignB, value_t, accum_t, epilogue_op_t, LdgAlignC, AllowRaggedTiles>,
m,
n,
k,
epilogue_op,
d_a,
d_b,
d_c,
stream,
debug_synchronous);
}
}
} // namespace gemm
} // namespace cutlass

View File

@ -1,653 +0,0 @@
/******************************************************************************
* Copyright (c) 2011-2017, NVIDIA CORPORATION. All rights reserved.
*
* Redistribution and use in source and binary forms, with or without
* modification, are not permitted.
*
* THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
* ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
* WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
* DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE FOR ANY
* DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
* (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
* LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
* ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
* (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
* SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
*
******************************************************************************/
#pragma once
/**
* \file
* Architecture-specific GEMM block_task policies
*/
#include <stdint.h>
#include "../util/util.h"
#include "block_task.h"
#include "grid_raster.h"
namespace cutlass {
namespace gemm {
/******************************************************************************
* tiling_strategy
******************************************************************************/
/**
* Enumeration of tile-sizing granularities
*/
struct tiling_strategy : printable_t
{
/// \brief Enumerants
enum kind_t
{
Unknown,
Small,
Medium,
Large,
Tall,
Wide,
Huge,
};
/// Enumerant value
kind_t kind;
/// Default constructor
tiling_strategy() : kind(Unknown) {}
/// Copy constructor
tiling_strategy(const kind_t &other_kind) : kind(other_kind) {}
/// Cast to kind_t
operator kind_t() const { return kind; }
/// Returns the instance as a string
__host__ __device__ inline
char const* to_string() const
{
switch (kind)
{
case Small: return "small";
case Medium: return "medium";
case Large: return "large";
case Tall: return "tall";
case Wide: return "wide";
case Huge: return "huge";
case Unknown:
default: return "unknown";
}
}
/// Insert the formatted instance into the output stream
void print(std::ostream& out) const { out << to_string(); }
};
/******************************************************************************
* GEMM
******************************************************************************/
/**
* GEMM task policy specialization for sgemm
*/
template <
typename value_t,
typename accum_t,
matrix_transform_t::kind_t TransformA, ///< Transformation op for matrix A
matrix_transform_t::kind_t TransformB, ///< Transformation op for matrix B
tiling_strategy::kind_t TilingStrategy> ///< Tile-sizing classification
struct gemm_policy;
/******************************************************************************
* SGEMM
******************************************************************************/
/**
* GEMM task policy specialization for small sgemm
*/
template <
matrix_transform_t::kind_t TransformA, ///< Transformation op for matrix A
matrix_transform_t::kind_t TransformB> ///< Transformation op for matrix B
struct gemm_policy<float, float, TransformA, TransformB, tiling_strategy::Small> :
block_task_policy<
16, // _BlockItemsY
16, // _BlockItemsX
16, // _BlockItemsK
2, // _ThreadItemsY
2, // _ThreadItemsX
false, // _UseDoubleScratchTiles
grid_raster_strategy::Default> // _RasterStrategy
{};
/**
* GEMM task policy specialization for medium sgemm
*/
template <
matrix_transform_t::kind_t TransformA, ///< Transformation op for matrix A
matrix_transform_t::kind_t TransformB> ///< Transformation op for matrix B
struct gemm_policy<float, float, TransformA, TransformB, tiling_strategy::Medium> :
block_task_policy<
32, // _BlockItemsY
32, // _BlockItemsX
8, // _BlockItemsK
4, // _ThreadItemsY
4, // _ThreadItemsX
false, // _UseDoubleScratchTiles
grid_raster_strategy::Default> // _RasterStrategy
{};
/**
* GEMM task policy specialization for large sgemm
*/
template <
matrix_transform_t::kind_t TransformA, ///< Transformation op for matrix A
matrix_transform_t::kind_t TransformB> ///< Transformation op for matrix B
struct gemm_policy<float, float, TransformA, TransformB, tiling_strategy::Large> :
block_task_policy<
64, // _BlockItemsY
64, // _BlockItemsX
8, // _BlockItemsK
8, // _ThreadItemsY
8, // _ThreadItemsX
false, // _UseDoubleScratchTiles
grid_raster_strategy::Default> // _RasterStrategy
{};
/**
* GEMM task policy specialization for tall sgemm
*/
template <
matrix_transform_t::kind_t TransformA, ///< Transformation op for matrix A
matrix_transform_t::kind_t TransformB> ///< Transformation op for matrix B
struct gemm_policy<float, float, TransformA, TransformB, tiling_strategy::Tall> :
block_task_policy<
128, // _BlockItemsY
32, // _BlockItemsX
8, // _BlockItemsK
8, // _ThreadItemsY
4, // _ThreadItemsX
false, // _UseDoubleScratchTiles
grid_raster_strategy::Default> // _RasterStrategy
{};
/**
* GEMM task policy specialization for wide sgemm
*/
template <
matrix_transform_t::kind_t TransformA, ///< Transformation op for matrix A
matrix_transform_t::kind_t TransformB> ///< Transformation op for matrix B
struct gemm_policy<float, float, TransformA, TransformB, tiling_strategy::Wide> :
block_task_policy<
32, // _BlockItemsY
128, // _BlockItemsX
8, // _BlockItemsK
4, // _ThreadItemsY
8, // _ThreadItemsX
false, // _UseDoubleScratchTiles
grid_raster_strategy::Default> // _RasterStrategy
{};
/**
* GEMM task policy specialization for huge sgemm
*/
template <
matrix_transform_t::kind_t TransformA, ///< Transformation op for matrix A
matrix_transform_t::kind_t TransformB> ///< Transformation op for matrix B
struct gemm_policy<float, float, TransformA, TransformB, tiling_strategy::Huge> :
block_task_policy<
128, // _BlockItemsY
128, // _BlockItemsX
8, // _BlockItemsK
8, // _ThreadItemsY
8, // _ThreadItemsX
false, // _UseDoubleScratchTiles
grid_raster_strategy::Default> // _RasterStrategy
{};
/******************************************************************************
* DGEMM
******************************************************************************/
/**
* GEMM task policy specialization for small dgemm
*/
template <
matrix_transform_t::kind_t TransformA, ///< Transformation op for matrix A
matrix_transform_t::kind_t TransformB> ///< Transformation op for matrix B
struct gemm_policy<double, double, TransformA, TransformB, tiling_strategy::Small> :
block_task_policy<
16, // _BlockItemsY
16, // _BlockItemsX
16, // _BlockItemsK
2, // _ThreadItemsY
2, // _ThreadItemsX
false, // _UseDoubleScratchTiles
grid_raster_strategy::Default> // _RasterStrategy
{};
/**
* GEMM task policy specialization for medium dgemm
*/
template <
matrix_transform_t::kind_t TransformA, ///< Transformation op for matrix A
matrix_transform_t::kind_t TransformB> ///< Transformation op for matrix B
struct gemm_policy<double, double, TransformA, TransformB, tiling_strategy::Medium> :
block_task_policy<
32, // _BlockItemsY
32, // _BlockItemsX
16, // _BlockItemsK
4, // _ThreadItemsY
4, // _ThreadItemsX
false, // _UseDoubleScratchTiles
grid_raster_strategy::Default> // _RasterStrategy
{};
/**
* GEMM task policy specialization for large dgemm
*/
template <
matrix_transform_t::kind_t TransformA, ///< Transformation op for matrix A
matrix_transform_t::kind_t TransformB> ///< Transformation op for matrix B
struct gemm_policy<double, double, TransformA, TransformB, tiling_strategy::Large> :
block_task_policy<
64, // _BlockItemsY
64, // _BlockItemsX
8, // _BlockItemsK
4, // _ThreadItemsY
4, // _ThreadItemsX
false, // _UseDoubleScratchTiles
grid_raster_strategy::Default> // _RasterStrategy
{};
/**
* GEMM task policy specialization for tall dgemm
*/
template <
matrix_transform_t::kind_t TransformA, ///< Transformation op for matrix A
matrix_transform_t::kind_t TransformB> ///< Transformation op for matrix B
struct gemm_policy<double, double, TransformA, TransformB, tiling_strategy::Tall> :
block_task_policy<
128, // _BlockItemsY
32, // _BlockItemsX
8, // _BlockItemsK
8, // _ThreadItemsY
4, // _ThreadItemsX
false, // _UseDoubleScratchTiles
grid_raster_strategy::Default> // _RasterStrategy
{};
/**
* GEMM task policy specialization for wide dgemm
*/
template <
matrix_transform_t::kind_t TransformA, ///< Transformation op for matrix A
matrix_transform_t::kind_t TransformB> ///< Transformation op for matrix B
struct gemm_policy<double, double, TransformA, TransformB, tiling_strategy::Wide> :
block_task_policy<
32, // _BlockItemsY
128, // _BlockItemsX
8, // _BlockItemsK
4, // _ThreadItemsY
8, // _ThreadItemsX
false, // _UseDoubleScratchTiles
grid_raster_strategy::Default> // _RasterStrategy
{};
/**
* GEMM task policy specialization for huge dgemm
*/
template <
matrix_transform_t::kind_t TransformA, ///< Transformation op for matrix A
matrix_transform_t::kind_t TransformB> ///< Transformation op for matrix B
struct gemm_policy<double, double, TransformA, TransformB, tiling_strategy::Huge> :
block_task_policy<
64, // _BlockItemsY
128, // _BlockItemsX
8, // _BlockItemsK
8, // _ThreadItemsY
8, // _ThreadItemsX
false, // _UseDoubleScratchTiles
grid_raster_strategy::Default> // _RasterStrategy
{};
/******************************************************************************
* HGEMM
******************************************************************************/
/**
* GEMM task policy specialization for small hgemm
*/
template <
matrix_transform_t::kind_t TransformA, ///< Transformation op for matrix A
matrix_transform_t::kind_t TransformB> ///< Transformation op for matrix B
struct gemm_policy<__half, __half, TransformA, TransformB, tiling_strategy::Small> :
block_task_policy<
32, // _BlockItemsY
32, // _BlockItemsX
8, // _BlockItemsK
4, // _ThreadItemsY
4, // _ThreadItemsX
false, // _UseDoubleScratchTiles
grid_raster_strategy::Default> // _RasterStrategy
{};
/**
* GEMM task policy specialization for medium hgemm
*/
template <
matrix_transform_t::kind_t TransformA, ///< Transformation op for matrix A
matrix_transform_t::kind_t TransformB> ///< Transformation op for matrix B
struct gemm_policy<__half, __half, TransformA, TransformB, tiling_strategy::Medium> :
block_task_policy<
32, // _BlockItemsY
32, // _BlockItemsX
16, // _BlockItemsK
8, // _ThreadItemsY
4, // _ThreadItemsX
false, // _UseDoubleScratchTiles
grid_raster_strategy::Default> // _RasterStrategy
{};
/**
* GEMM task policy specialization for large hgemm
*/
template <
matrix_transform_t::kind_t TransformA, ///< Transformation op for matrix A
matrix_transform_t::kind_t TransformB> ///< Transformation op for matrix B
struct gemm_policy<__half, __half, TransformA, TransformB, tiling_strategy::Large> :
block_task_policy<
64, // _BlockItemsY
64, // _BlockItemsX
8, // _BlockItemsK
16, // _ThreadItemsY
8, // _ThreadItemsX
false, // _UseDoubleScratchTiles
grid_raster_strategy::Default> // _RasterStrategy
{};
/**
* GEMM task policy specialization for tall hgemm
*/
template <
matrix_transform_t::kind_t TransformA, ///< Transformation op for matrix A
matrix_transform_t::kind_t TransformB> ///< Transformation op for matrix B
struct gemm_policy<__half, __half, TransformA, TransformB, tiling_strategy::Tall> :
block_task_policy<
128, // _BlockItemsY
32, // _BlockItemsX
8, // _BlockItemsK
16, // _ThreadItemsY
4, // _ThreadItemsX
false, // _UseDoubleScratchTiles
grid_raster_strategy::Default> // _RasterStrategy
{};
/**
* GEMM task policy specialization for wide hgemm
*/
template <
matrix_transform_t::kind_t TransformA, ///< Transformation op for matrix A
matrix_transform_t::kind_t TransformB> ///< Transformation op for matrix B
struct gemm_policy<__half, __half, TransformA, TransformB, tiling_strategy::Wide> :
block_task_policy<
32, // _BlockItemsY
128, // _BlockItemsX
8, // _BlockItemsK
8, // _ThreadItemsY
8, // _ThreadItemsX
false, // _UseDoubleScratchTiles
grid_raster_strategy::Default> // _RasterStrategy
{};
/**
* GEMM task policy specialization for huge hgemm
*/
template <
matrix_transform_t::kind_t TransformA, ///< Transformation op for matrix A
matrix_transform_t::kind_t TransformB> ///< Transformation op for matrix B
struct gemm_policy<__half, __half, TransformA, TransformB, tiling_strategy::Huge> :
block_task_policy<
128, // _BlockItemsY
128, // _BlockItemsX
8, // _BlockItemsK
16, // _ThreadItemsY
8, // _ThreadItemsX
false, // _UseDoubleScratchTiles
grid_raster_strategy::Default> // _RasterStrategy
{};
/******************************************************************************
* IGEMM
******************************************************************************/
/**
* GEMM task policy specialization for small igemm
*/
template <
matrix_transform_t::kind_t TransformA, ///< Transformation op for matrix A
matrix_transform_t::kind_t TransformB> ///< Transformation op for matrix B
struct gemm_policy<int8_t, int32_t, TransformA, TransformB, tiling_strategy::Small> :
block_task_policy<
16, // _BlockItemsY
32, // _BlockItemsX
32, // _BlockItemsK
4, // _ThreadItemsY
4, // _ThreadItemsX
false, // _UseDoubleScratchTiles
grid_raster_strategy::Default> // _RasterStrategy
{};
/**
* GEMM task policy specialization for medium igemm
*/
template <
matrix_transform_t::kind_t TransformA, ///< Transformation op for matrix A
matrix_transform_t::kind_t TransformB> ///< Transformation op for matrix B
struct gemm_policy<int8_t, int32_t, TransformA, TransformB, tiling_strategy::Medium> :
block_task_policy<
32, // _BlockItemsY
32, // _BlockItemsX
32, // _BlockItemsK
4, // _ThreadItemsY
4, // _ThreadItemsX
false, // _UseDoubleScratchTiles
grid_raster_strategy::Default> // _RasterStrategy
{};
/**
* GEMM task policy specialization for large igemm
*/
template <
matrix_transform_t::kind_t TransformA, ///< Transformation op for matrix A
matrix_transform_t::kind_t TransformB> ///< Transformation op for matrix B
struct gemm_policy<int8_t, int32_t, TransformA, TransformB, tiling_strategy::Large> :
block_task_policy<
64, // _BlockItemsY
64, // _BlockItemsX
32, // _BlockItemsK
8, // _ThreadItemsY
4, // _ThreadItemsX
false, // _UseDoubleScratchTiles
grid_raster_strategy::Default> // _RasterStrategy
{};
/**
* GEMM task policy specialization for large igemm
*/
template <
matrix_transform_t::kind_t TransformA, ///< Transformation op for matrix A
matrix_transform_t::kind_t TransformB> ///< Transformation op for matrix B
struct gemm_policy<int8_t, int32_t, TransformA, TransformB, tiling_strategy::Tall> :
block_task_policy<
128, // _BlockItemsY
64, // _BlockItemsX
64, // _BlockItemsK
8, // _ThreadItemsY
4, // _ThreadItemsX
false, // _UseDoubleScratchTiles
grid_raster_strategy::Default> // _RasterStrategy
{};
/**
* GEMM task policy specialization for large igemm
*/
template <
matrix_transform_t::kind_t TransformA, ///< Transformation op for matrix A
matrix_transform_t::kind_t TransformB> ///< Transformation op for matrix B
struct gemm_policy<int8_t, int32_t, TransformA, TransformB, tiling_strategy::Wide> :
block_task_policy<
64, // _BlockItemsY
128, // _BlockItemsX
64, // _BlockItemsK
4, // _ThreadItemsY
8, // _ThreadItemsX
false, // _UseDoubleScratchTiles
grid_raster_strategy::Default> // _RasterStrategy
{};
/**
* GEMM task policy specialization for huge igemm
*/
template <
matrix_transform_t::kind_t TransformA, ///< Transformation op for matrix A
matrix_transform_t::kind_t TransformB> ///< Transformation op for matrix B
struct gemm_policy<int8_t, int32_t, TransformA, TransformB, tiling_strategy::Huge> :
block_task_policy<
128, // _BlockItemsY
128, // _BlockItemsX
32, // _BlockItemsK
8, // _ThreadItemsY
8, // _ThreadItemsX
false, // _UseDoubleScratchTiles
grid_raster_strategy::Default> // _RasterStrategy
{};
/******************************************************************************
* WMMA GEMM
******************************************************************************/
// WMMA is a preview feature in CUDA. Conditionally enable wmma_gemm policies.
#if defined(WMMA)
template <
matrix_transform_t::kind_t TransformA, ///< Transformation op for matrix A
matrix_transform_t::kind_t TransformB> ///< Transformation op for matrix B
struct gemm_policy<half, float, TransformA, TransformB, tiling_strategy::Small> :
gemm::block_task_wmma_policy<
16, // _BlockItemsY
16, // _BlockItemsX
16, // _BlockItemsK
16, // _WarpItemsY
16, // _WarpItemsX
16, // _WmmaItemsY
16, // _WmmaItemsX
16, // _WmmaItemsK
false, // _UseDoubleScratchTiles
gemm::grid_raster_strategy::Default> // _RasterStrategy
{};
template <
matrix_transform_t::kind_t TransformA, ///< Transformation op for matrix A
matrix_transform_t::kind_t TransformB> ///< Transformation op for matrix B
struct gemm_policy<half, float, TransformA, TransformB, tiling_strategy::Medium> :
gemm::block_task_wmma_policy<
32, // _BlockItemsY
32, // _BlockItemsX
32, // _BlockItemsK
32, // _WarpItemsY
32, // _WarpItemsX
16, // _WmmaItemsY
16, // _WmmaItemsX
16, // _WmmaItemsK
false, // _UseDoubleScratchTiles
gemm::grid_raster_strategy::Default> // _RasterStrategy
{};
template <
matrix_transform_t::kind_t TransformA, ///< Transformation op for matrix A
matrix_transform_t::kind_t TransformB> ///< Transformation op for matrix B
struct gemm_policy< half, float, TransformA, TransformB, tiling_strategy::Large> :
gemm::block_task_wmma_policy<
64, // _BlockItemsY
64, // _BlockItemsX
32, // _BlockItemsK
32, // _WarpItemsY
64, // _WarpItemsX
16, // _WmmaItemsY
16, // _WmmaItemsX
16, // _WmmaItemsK
false, // _UseDoubleScratchTiles
gemm::grid_raster_strategy::Default> // _RasterStrategy
{};
template <
matrix_transform_t::kind_t TransformA, ///< Transformation op for matrix A
matrix_transform_t::kind_t TransformB> ///< Transformation op for matrix B
struct gemm_policy< half, float, TransformA, TransformB, tiling_strategy::Tall> :
gemm::block_task_wmma_policy<
128, // _BlockItemsY
64, // _BlockItemsX
64, // _BlockItemsK
32, // _WarpItemsY
64, // _WarpItemsX
16, // _WmmaItemsY
16, // _WmmaItemsX
16, // _WmmaItemsK
false, // _UseDoubleScratchTiles
gemm::grid_raster_strategy::Default> // _RasterStrategy
{};
template <
matrix_transform_t::kind_t TransformA, ///< Transformation op for matrix A
matrix_transform_t::kind_t TransformB> ///< Transformation op for matrix B
struct gemm_policy< half, float, TransformA, TransformB, tiling_strategy::Wide> :
gemm::block_task_wmma_policy<
64, // _BlockItemsY
128, // _BlockItemsX
64, // _BlockItemsK
32, // _WarpItemsY
64, // _WarpItemsX
16, // _WmmaItemsY
16, // _WmmaItemsX
16, // _WmmaItemsK
false, // _UseDoubleScratchTiles
gemm::grid_raster_strategy::Default> // _RasterStrategy
{};
template <
matrix_transform_t::kind_t TransformA, ///< Transformation op for matrix A
matrix_transform_t::kind_t TransformB> ///< Transformation op for matrix B
struct gemm_policy< half, float, TransformA, TransformB, tiling_strategy::Huge> :
gemm::block_task_wmma_policy<
128, // _BlockItemsY
128, // _BlockItemsX
64, // _BlockItemsK
32, // _WarpItemsY
64, // _WarpItemsX
16, // _WmmaItemsY
16, // _WmmaItemsX
16, // _WmmaItemsK
false, // _UseDoubleScratchTiles
gemm::grid_raster_strategy::Default> // _RasterStrategy
{};
#endif
} // namespace gemm
} // namespace cutlass

View File

@ -1,215 +0,0 @@
/******************************************************************************
* Copyright (c) 2011-2017, NVIDIA CORPORATION. All rights reserved.
*
* Redistribution and use in source and binary forms, with or without
* modification, are not permitted.
*
* THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
* ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
* WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
* DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE FOR ANY
* DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
* (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
* LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
* ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
* (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
* SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
*
******************************************************************************/
#pragma once
/**
* \file
* Abstraction for exposing architecture-specific "dot-product-accumulate"
* ISA operations
*/
#include <stdint.h>
#include "../util/util.h"
namespace cutlass {
namespace gemm {
/******************************************************************************
* dp_accummulate
******************************************************************************/
/**
* \brief Abstraction for exposing architecture-specific "dot-product-accumulate"
* ISA operations
*
* Given two K-component vectors a and b having type value_t[K] and an addend c
* of type accum_t, the "dot-product-accumulate" of type accum_t is computed
* as d = x[0]*y[0] + x[1]*y[1] + ... + x[K-1]*y[K-1] + c.
*
* We use the notation "dpK" to connote a K-component dot-product-accumulate.
* For example, "dp1" is a simple multiply-add.
*
* For given pairing of value_t and accum_t types, the corresponding
* dp_accummulate class will:
*
* - Define the member-type dp_vector_t as the appropriate K-component vector
* type needed to leverage architecture-specific "dot-product accumulate"
* ISA operations.
* - Implement the corresponding dot-product operation between two dp_vector_t
* inputs a and b.
*
*/
template <
typename value_t, ///< Component value type
typename accum_t> ///< Accumulator value type
struct dp_accummulate;
/// Default "dp1" dot-product-accumulate traits specialization for value_t->accum_t
template <
typename value_t, ///< Component value type
typename accum_t> ///< Accumulator value type
struct dp_accummulate
{
/// Single-component "dp1" dot-product vector type
typedef value_t dp_vector_t;
/// Compute "dp1" float->float
inline __device__
static void mad(
float &d,
const float &a,
const float &b,
const float &c)
{
asm volatile ( "fma.rn.f32 %0, %1, %2, %3;\n"
: "=f"(d) : "f"(a), "f"(b), "f"(c));
}
/// Compute "dp1" double->double
inline __device__
static void mad(
double &d,
const double &a,
const double &b,
const double &c)
{
asm volatile ("fma.rn.f64 %0, %1, %2, %3;\n"
: "=d"(d) : "d"(a), "d"(b), "d"(c));
}
/// Compute "dp1" int16_t->int32_t
inline __device__
static void mad(
int32_t &d,
const int16_t &a,
const int16_t &b,
const int32_t &c)
{
asm volatile ("mad.wide.s16 %0, %1, %2, %3;\n"
: "=r"(d) : "h"(a), "h"(b), "r"(c));
}
/// Compute "dp1" uint16_t->uint32_t
inline __device__
static void mad(
uint32_t &d,
const uint16_t &a,
const uint16_t &b,
const uint32_t &c)
{
asm volatile ("mad.wide.u16 %0, %1, %2, %3;\n"
: "=r"(d) : "h"(a), "h"(b), "r"(c));
}
/// Compute "dp1" int32_t->int32_t
inline __device__
static void mad(
int32_t &d,
const int32_t &a,
const int32_t &b,
const int32_t &c)
{
asm volatile ("mad.lo.s32 %0, %1, %2, %3;\n"
: "=r"(d) : "r"(a), "r"(b), "r"(c));
}
/// Compute "dp1" uint32_t->uint32_t
inline __device__
static void mad(
uint32_t &d,
const uint32_t &a,
const uint32_t &b,
const uint32_t &c)
{
asm volatile ("mad.lo.u32 %0, %1, %2, %3;\n"
: "=r"(d) : "r"(a), "r"(b), "r"(c));
}
};
#if (CUTLASS_ARCH >= 610) // Specializations only enabled for Pascal SM610+
/// "dp4" dot-product-accumulate traits specialization for int8_t->int32_t
template <>
struct dp_accummulate<
int8_t, ///< Component value type
int32_t> ///< Accumulator value type
{
/// Four-component signed "idp4"
typedef int32_t dp_vector_t;
/// Compute "dp4" int16_t->int32_t
inline __device__
static void mad(
int32_t &d,
const int32_t &a,
const int32_t &b,
const int32_t &c)
{
asm volatile ( "dp4a.s32.s32 %0, %1, %2, %3;\n"
: "=r"(d) : "r"(a), "r"(b), "r"(c));
}
};
/// "dp4" dot-product-accumulate traits specialization for uint8_t->uint32_t
template <>
struct dp_accummulate<
uint8_t, ///< Component value type
uint32_t> ///< Accumulator value type
{
/// Four-component unsigned "idp4"
typedef uint32_t dp_vector_t;
/// Compute "dp4" uint16_t->uint32_t
inline __device__
static void mad(
uint32_t &d,
const uint32_t &a,
const uint32_t &b,
const uint32_t &c)
{
asm volatile ( "dp4a.u32.u32 %0, %1, %2, %3;\n"
: "=r"(d) : "r"(a), "r"(b), "r"(c));
}
};
#endif // Specializations only enabled for Pascal SM610+
} // namespace gemm
} // namespace cutlass

View File

@ -1,96 +0,0 @@
/******************************************************************************
* Copyright (c) 2011-2017, NVIDIA CORPORATION. All rights reserved.
*
* Redistribution and use in source and binary forms, with or without
* modification, are not permitted.
*
* THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
* ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
* WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
* DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE FOR ANY
* DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
* (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
* LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
* ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
* (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
* SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
*
******************************************************************************/
#pragma once
/**
* \file
* Epilogue operation to compute final output
*/
namespace cutlass {
namespace gemm {
//// Used by GEMM to compute the final result C <= alpha * accumulator + beta * C
template <
typename accum_t,
typename output_t,
typename scalar_t
>
class blas_scaled_epilogue
{
public:
scalar_t alpha;
scalar_t beta;
inline __device__ __host__
blas_scaled_epilogue(
scalar_t alpha,
scalar_t beta)
:
alpha(alpha),
beta(beta)
{}
/// Epilogue operator
inline __device__ __host__
output_t operator()(
accum_t accumulator,
output_t c,
size_t idx) const
{
return output_t(alpha * scalar_t(accumulator) + beta * scalar_t(c));
}
/// Epilogue operator
inline __device__ __host__
output_t operator()(
accum_t accumulator,
size_t idx) const
{
return output_t(alpha * scalar_t(accumulator));
}
/**
* Configure epilogue as to whether the thread block is a secondary
* accumulator in an inter-block k-splitting scheme
*/
inline __device__
void set_secondary_accumulator()
{
beta = scalar_t(1);
}
/// Return whether the beta-scaled addend needs initialization
inline __device__
bool must_init_addend()
{
return (beta != scalar_t(0));
}
};
} // namespace gemm
} // namespace cutlass

View File

@ -1,428 +0,0 @@
/******************************************************************************
* Copyright (c) 2011-2017, NVIDIA CORPORATION. All rights reserved.
*
* Redistribution and use in source and binary forms, with or without
* modification, are not permitted.
*
* THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
* ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
* WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
* DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE FOR ANY
* DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
* (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
* LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
* ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
* (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
* SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
*
******************************************************************************/
#pragma once
/**
* \file
* Abstraction for enumerating \p block_task within an input matrix
*/
#include <stdint.h>
#include "../util/util.h"
namespace cutlass {
namespace gemm {
/******************************************************************************
* grid_raster_strategy
******************************************************************************/
/**
* \brief Strategies for enumerating \p block_task within an input matrix
*/
struct grid_raster_strategy
{
/// \brief Enumerants
enum kind_t
{
/**
* Default \p block_task assignment (currently ColumnMajor for N*,
* RowMajor for TT, and TiledCohort for TN)
*/
Default,
/**
* Column-major \p block_task assignment
*/
ColumnMajor,
/**
* Row-major \p block_task assignment
*/
RowMajor,
/**
* Two-level \p block_task assignment (both column-major)
*/
TiledCohort,
};
};
/******************************************************************************
* grid_raster
******************************************************************************/
/**
* \brief Abstraction for enumerating \p block_task within an input matrix
*
* NB: This generic class is not directly constructible. Algorithm-specific
* template specializations will provide the API functionality prescribed here.
*/
template <
int BlockItemsY, ///< Height in rows of a block-wide tile in matrix C
int BlockItemsX, ///< Width in columns of a block-wide tile in matrix C
matrix_transform_t::kind_t TransformA, ///< View transform enumerant for matrix A
matrix_transform_t::kind_t TransformB, ///< View transform enumerant for matrix B
grid_raster_strategy::kind_t RasterStrategy> ///< Strategy for enumerating \p block_task within an input matrix
struct grid_raster
{
//-------------------------------------------------------------------------
// Device API
//-------------------------------------------------------------------------
/// Thread block's base item coordinates (x, y) in matrix C
int2 block_item_coords;
/// Constructor
grid_raster();
/// Whether the thread block base coordinates are out-of-bounds for an m*n matrix C
bool is_block_oob(int m, int n);
//-------------------------------------------------------------------------
// Grid launch API
//-------------------------------------------------------------------------
/// Compute the kernel grid extents (in thread blocks) for consuming an m*n matrix C
static dim3 grid_dims(int m, int n);
};
/******************************************************************************
* grid_raster (ColumnMajor specialization)
******************************************************************************/
/**
* \brief Abstraction for enumerating \p block_task within an input matrix
* (ColumnMajor specialization)
*
* Maps thread blocksin column-major fashion
*/
template <
int BlockItemsY, ///< Height in rows of a block-wide tile in matrix C
int BlockItemsX, ///< Width in columns of a block-wide tile in matrix C
matrix_transform_t::kind_t TransformA, ///< View transform enumerant for matrix A
matrix_transform_t::kind_t TransformB> ///< View transform enumerant for matrix B
struct grid_raster<
BlockItemsY,
BlockItemsX,
TransformA,
TransformB,
grid_raster_strategy::ColumnMajor> ///< Strategy for enumerating \p block_task within an input matrix
{
//-------------------------------------------------------------------------
// Device API
//-------------------------------------------------------------------------
/// Thread block's base item coordinates (x, y) in matrix C
int2 block_item_coords;
/// Constructor
inline __device__
grid_raster()
{
// blockDim.x is the fastest changing grid dim on current architectures
block_item_coords = make_int2(
BlockItemsX * blockIdx.y,
BlockItemsY * blockIdx.x);
}
/// Whether the base \p block_item_coords are out-of-bounds for an m*n matrix C
inline __device__
bool is_block_oob(int m, int n)
{
// ColumnMajor never rasterizes fully out-of-bounds thread blocks
return false;
}
//-------------------------------------------------------------------------
// Grid launch API
//-------------------------------------------------------------------------
/// Compute the kernel grid extents (in thread blocks) for consuming an m*n matrix C
inline __host__ __device__
static dim3 grid_dims(int m, int n)
{
// blockDim.x is the fastest changing grid dim on current architectures
return dim3(
(m + BlockItemsY - 1) / BlockItemsY,
(n + BlockItemsX - 1) / BlockItemsX);
}
};
/******************************************************************************
* grid_raster (RowMajor specialization)
******************************************************************************/
/**
* \brief Abstraction for enumerating \p block_task within an input matrix
* (RowMajor specialization)
*
* Enumerates \p block_task in row-major fashion
*/
template <
int BlockItemsY, ///< Height in rows of a block-wide tile in matrix C
int BlockItemsX, ///< Width in columns of a block-wide tile in matrix C
matrix_transform_t::kind_t TransformA, ///< View transform enumerant for matrix A
matrix_transform_t::kind_t TransformB> ///< View transform enumerant for matrix B
struct grid_raster<
BlockItemsY,
BlockItemsX,
TransformA,
TransformB,
grid_raster_strategy::RowMajor> ///< Strategy for enumerating \p block_task within an input matrix
{
//-------------------------------------------------------------------------
// Device API
//-------------------------------------------------------------------------
/// Thread block's base item coordinates (x, y) in matrix C
int2 block_item_coords;
/// Constructor
inline __device__
grid_raster()
{
// blockDim.x is the fastest changing grid dim on current architectures
block_item_coords = make_int2(
BlockItemsX * blockIdx.x,
BlockItemsY * blockIdx.y);
}
/// Whether the base \p block_item_coords are out-of-bounds for an m*n matrix C
inline __device__
bool is_block_oob(int m, int n)
{
// RowMajor never rasterizes fully out-of-bounds thread blocks
return false;
}
//-------------------------------------------------------------------------
// Grid launch API
//-------------------------------------------------------------------------
/// Compute the kernel grid extents (in thread blocks) for consuming an m*n matrix C
inline __host__ __device__
static dim3 grid_dims(int m, int n)
{
// blockDim.x is the fastest changing grid dim on current architectures
return dim3(
(n + BlockItemsX - 1) / BlockItemsX,
(m + BlockItemsY - 1) / BlockItemsY);
}
};
/******************************************************************************
* grid_raster (TiledCohort specialization)
******************************************************************************/
/**
* \brief Abstraction for enumerating \p block_task within an input matrix
* (TiledCohort specialization)
*
* Enumerates \p block_task in column-major fashion across "cohort" tiles (where
* cohorts are CohortBlocksY high and CohortBlocksX wide), and enumerates cohorts
* across the matrix in column-major fashion.
*
* Grid layout:
* - gridDim.y is the height of the grid in cohorts
* - gridDim.x is the width of the grid in cohorts multiplied by the number of
* thread blocks per cohort
*/
template <
int BlockItemsY, ///< Height in rows of a block-wide tile in matrix C
int BlockItemsX, ///< Width in columns of a block-wide tile in matrix C
matrix_transform_t::kind_t TransformA, ///< View transform enumerant for matrix A
matrix_transform_t::kind_t TransformB> ///< View transform enumerant for matrix B
struct grid_raster<
BlockItemsY,
BlockItemsX,
TransformA,
TransformB,
grid_raster_strategy::TiledCohort> ///< Strategy for enumerating \p block_task within an input matrix
{
enum
{
/// Height in thread blocks of a grid rasterization cohort
CohortBlocksY = 2,
/// Width in thread blocks of a grid rasterization cohort
CohortBlocksX = 2,
/// Number of thread blocks per cohort
BlocksPerCohort = CohortBlocksY * CohortBlocksX,
/// Height in items of a grid rasterization cohort
CohortItemsY = CohortBlocksY * BlockItemsY,
/// Width in items of a grid rasterization cohort
CohortItemsX = CohortBlocksX * BlockItemsX,
};
//-------------------------------------------------------------------------
// Device API
//-------------------------------------------------------------------------
/// Thread block's base item coordinates (x, y) in matrix C
int2 block_item_coords;
/// Constructor
inline __device__
grid_raster()
{
int block_idx_cohort = blockIdx.x % BlocksPerCohort;
int2 cohort_coords_grid = make_int2(
blockIdx.x / BlocksPerCohort,
blockIdx.y);
// Cohort is rastered in column-major order
int2 block_coords_cohort = make_int2(
block_idx_cohort / CohortBlocksY,
block_idx_cohort % CohortBlocksY);
block_item_coords = make_int2(
((cohort_coords_grid.x * CohortBlocksX) + block_coords_cohort.x) * BlockItemsX,
((cohort_coords_grid.y * CohortBlocksY) + block_coords_cohort.y) * BlockItemsY);
}
/// Whether the base \p block_item_coords are out-of-bounds for an m*n matrix C
inline __device__
bool is_block_oob(int m, int n)
{
/// thread blocks within the cohort may be fully out-of-bounds
return (block_item_coords.x >= n) || (block_item_coords.y >= m);
}
//-------------------------------------------------------------------------
// Grid launch API
//-------------------------------------------------------------------------
/// Compute the kernel grid extents (in thread blocks) for consuming an m*n matrix C
inline __host__ __device__
static dim3 grid_dims(int m, int n)
{
// Extents of C matrix in cohorts
int2 grid_cohort_dims = make_int2(
(n + CohortItemsX - 1) / CohortItemsX,
(m + CohortItemsY - 1) / CohortItemsY);
return dim3(
grid_cohort_dims.x * BlocksPerCohort, // gridDim.x is width of grid in cohorts * size of cohort in blocks
grid_cohort_dims.y, // gridDim.y is height of grid in cohorts
1); // gridDim.z is reserved for optional k-splitting
}
};
/******************************************************************************
* grid_raster (Default specializations)
******************************************************************************/
/**
* \brief Abstraction for enumerating \p block_task within an input matrix
* (Default N* specialization)
*
* Maps thread blocksin column-major fashion
*/
template <
int BlockItemsY, ///< Height in rows of a block-wide tile in matrix C
int BlockItemsX, ///< Width in columns of a block-wide tile in matrix C
matrix_transform_t::kind_t TransformB> ///< View transform enumerant for matrix B
struct grid_raster<
BlockItemsY,
BlockItemsX,
matrix_transform_t::NonTranspose, ///< View transform enumerant for matrix A
TransformB,
grid_raster_strategy::Default> ///< Strategy for enumerating \p block_task within an input matrix
:
grid_raster<
BlockItemsY,
BlockItemsX,
matrix_transform_t::NonTranspose,
TransformB,
grid_raster_strategy::ColumnMajor>
{};
/**
* \brief Abstraction for enumerating \p block_task within an input matrix
* (Default TT specialization)
*
* Maps thread blocksin row-major fashion
*/
template <
int BlockItemsY, ///< Height in rows of a block-wide tile in matrix C
int BlockItemsX> ///< Width in columns of a block-wide tile in matrix C
struct grid_raster<
BlockItemsY,
BlockItemsX,
matrix_transform_t::Transpose, ///< View transform enumerant for matrix A
matrix_transform_t::Transpose, ///< View transform enumerant for matrix B
grid_raster_strategy::Default> ///< Strategy for enumerating \p block_task within an input matrix
:
grid_raster<
BlockItemsY,
BlockItemsX,
matrix_transform_t::Transpose,
matrix_transform_t::Transpose,
grid_raster_strategy::RowMajor>
{};
/**
* \brief Abstraction for enumerating \p block_task within an input matrix
* (Default TN specialization)
*
* Maps thread blocksin blocked cohorts
*/
template <
int BlockItemsY, ///< Height in rows of a block-wide tile in matrix C
int BlockItemsX> ///< Width in columns of a block-wide tile in matrix C
struct grid_raster<
BlockItemsY,
BlockItemsX,
matrix_transform_t::Transpose, ///< View transform enumerant for matrix A
matrix_transform_t::NonTranspose, ///< View transform enumerant for matrix B
grid_raster_strategy::Default> ///< Strategy for enumerating \p block_task within an input matrix
:
grid_raster<
BlockItemsY,
BlockItemsX,
matrix_transform_t::Transpose,
matrix_transform_t::NonTranspose,
grid_raster_strategy::TiledCohort>
{};
} // namespace gemm
} // namespace cutlass

View File

@ -1,302 +0,0 @@
/******************************************************************************
* Copyright (c) 2011-2017, NVIDIA CORPORATION. All rights reserved.
*
* Redistribution and use in source and binary forms, with or without
* modification, are not permitted.
*
* THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
* ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
* WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
* DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE FOR ANY
* DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
* (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
* LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
* ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
* (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
* SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
*
******************************************************************************/
#pragma once
/**
* \file
* Abstraction for coordinating inter-block k-splitting
*/
#include <stdint.h>
#include "../util/util.h"
namespace cutlass {
namespace gemm {
/******************************************************************************
* Storage and initialization
******************************************************************************/
enum
{
NumFlagsSplitK = 4096
};
/**
* Global K-split semaphore flags
*
* TODO: use demand-allocated storage to provide copies for concurrent streams
*/
__device__ int d_flags_split_k[NumFlagsSplitK];
/**
* Preparation kernel for zero-initializing semaphore flags
*/
__global__ void prepare_kernel(int *d_flags_split_k)
{
int tid = (blockIdx.x * blockDim.x) + threadIdx.x;
if (tid < NumFlagsSplitK)
d_flags_split_k[tid] = 0;
}
/******************************************************************************
* k_split_control
******************************************************************************/
/**
* \brief Abstraction for coordinating inter-block k-splitting
*/
struct k_split_control
{
/// Extent of a thread block's partition along the GEMM K-axis
int split_k;
/// Whether or not to use a semaphore for inter-block k-splitting.
bool use_semaphore;
/// Pointer to semaphore
int *d_flags;
//-------------------------------------------------------------------------
// Device API
//-------------------------------------------------------------------------
/**
* Return the thread block's starting coordinate (k) within the
* multiplicand matrices
*/
inline __device__
int block_begin_item_k()
{
return blockIdx.z * split_k;
}
/**
* Return the thread block's ending coordinate (k) within the multiplicand
* matrices (one-past)
*/
inline __device__
int block_end_item_k(int dim_k)
{
int next_start_k = block_begin_item_k() + split_k;
return __NV_STD_MIN(next_start_k, dim_k);
}
/**
* Whether the thread block is a secondary accumulator in an inter-block
* k-splitting scheme
*/
inline __device__
bool is_secondary_accumulator()
{
return (blockIdx.z > 0);
}
/**
* Wait for predecessor thread block(s) to produce the exclusive
* partial-sums for this block-wide tile
*/
inline __device__
void wait()
{
// Wait on semaphore
if ((use_semaphore) && (blockIdx.z > 0))
{
if (threadIdx.x == 0)
{
int bid = (blockIdx.y * gridDim.x) + blockIdx.x;
int hash = bid % NumFlagsSplitK;
int found;
int looking = blockIdx.z;
while (true)
{
asm volatile ("ld.global.cg.u32 %0, [%1];\n" : "=r"(found) : "l"(d_flags + hash));
if (found == looking)
break;
/// Fence to keep load from being hoisted from the loop
__syncwarp(0x00000001);
}
}
__syncthreads();
}
}
/**
* Signal the successor thread_block(s) that the inclusive partial-sums
* from this block-wide tile are available
*/
inline __device__
void signal()
{
if (use_semaphore)
{
__syncthreads();
if (threadIdx.x == 0)
{
int bid = (blockIdx.y * gridDim.x) + blockIdx.x;
int hash = bid % NumFlagsSplitK;
int val = blockIdx.z + 1;
asm volatile ("st.global.cg.u32 [%0], %1;\n" : : "l"(d_flags + hash), "r"(val));
}
}
}
//-------------------------------------------------------------------------
// Grid launch API
//-------------------------------------------------------------------------
/**
* Constructor
*/
inline
k_split_control(
int *d_flags,
int sm_count,
int max_sm_occupancy,
int dim_k,
int block_tile_items_k,
dim3 block_dims,
dim3 &grid_dims) ///< [in,out]
:
d_flags(d_flags),
split_k(dim_k)
{
// Compute wave efficiency
float wave_efficiency = get_wave_efficiency(
sm_count,
max_sm_occupancy,
block_dims,
grid_dims);
// Update split-k if wave efficiency is less than some threshold
if (wave_efficiency < 0.9)
{
int num_threadblocks = grid_dims.x * grid_dims.y * grid_dims.z;
// Ideal number of thread blocks in grid
int ideal_threadblocks = lcm(sm_count, num_threadblocks);
// Desired number of partitions to split K-axis into
int num_partitions = ideal_threadblocks / num_threadblocks;
// Compute new k-split share
int new_split_k = (dim_k + num_partitions - 1) / num_partitions;
// Round split_k share to the nearest block_task_policy_t::BlockItemsK
new_split_k = round_nearest(new_split_k, block_tile_items_k);
// Recompute k-splitting factor with new_split_k
num_partitions = (dim_k + new_split_k - 1) / new_split_k;
// Update grid dims and k if we meet the minimum number of iterations worth the overhead of splitting
int min_iterations_k = 8;
if (((new_split_k / block_tile_items_k) > min_iterations_k) && // We're going to go through at least this many k iterations
(sm_count * max_sm_occupancy < NumFlagsSplitK)) // We have enough semaphore flags allocated
{
grid_dims.z = num_partitions;
split_k = new_split_k;
}
}
use_semaphore = (grid_dims.z > 1);
}
/**
* Initializer
*/
cudaError_t prepare(
cudaStream_t stream, ///< CUDA stream to launch kernels within. Default is stream<sub>0</sub>.
bool debug_synchronous) ///< Whether or not to synchronize the stream after every kernel launch to check for errors. Also causes launch configurations to be printed to the console if DEBUG is defined. Default is \p false.
{
cudaError error = cudaSuccess;
if (use_semaphore)
{
int block_threads = 128;
int grid_dims = (NumFlagsSplitK + block_threads - 1) / block_threads;
prepare_kernel<<<grid_dims, block_threads, 0, stream>>>(d_flags);
// Check for failure to launch
if (CUDA_PERROR_DEBUG(error = cudaPeekAtLastError()))
return error;
// Sync the stream if specified to flush runtime errors
if (debug_synchronous && (CUDA_PERROR_DEBUG(error = cudaStreamSynchronize(stream))))
return error;
}
return error;
}
/**
* Compute the efficiency of dispatch wave quantization
*/
float get_wave_efficiency(
int sm_count,
int max_sm_occupancy,
dim3 block_dims,
dim3 grid_dims)
{
// Heuristic for how many warps are needed to saturate an SM for a given
// multiply-accumulate genre. (NB: We could make this more rigorous by
// specializing on data types and SM width)
int saturating_warps_per_sm = 16;
int num_threadblocks = grid_dims.x * grid_dims.y * grid_dims.z;
int threads_per_threadblock = block_dims.x * block_dims.y;
int warps_per_threadblock = threads_per_threadblock / 32;
int saturating_threadblocks_per_sm = (saturating_warps_per_sm + warps_per_threadblock - 1) / warps_per_threadblock;
int saturating_residency = sm_count * saturating_threadblocks_per_sm;
int full_waves = num_threadblocks / saturating_residency;
int remainder_threadblocks = num_threadblocks % saturating_residency;
int total_waves = (remainder_threadblocks == 0) ? full_waves : full_waves + 1;
float last_wave_saturating_efficiency = float(remainder_threadblocks) / saturating_residency;
return (float(full_waves) + last_wave_saturating_efficiency) / total_waves;
}
};
} // namespace gemm
} // namespace cutlass

View File

@ -1,461 +0,0 @@
/******************************************************************************
* Copyright (c) 2011-2017, NVIDIA CORPORATION. All rights reserved.
*
* Redistribution and use in source and binary forms, with or without
* modification, are not permitted.
*
* THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
* ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
* WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
* DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE FOR ANY
* DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
* (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
* LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
* ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
* (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
* SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
*
******************************************************************************/
#pragma once
/**
* \file
* Thread-level multiply-accumulate abstraction
*/
#include "../util/util.h"
#include "dp_accummulate.h"
namespace cutlass {
namespace gemm {
/******************************************************************************
* thread_accumulator (generic specialization)
******************************************************************************/
/**
* \brief Thread-level multiply-accumulate abstraction (generic specialization)
*
* The thread_accumulator class maintains a MxN tile of accumulators in
* registers to which MxNxK matrix products of two thread tiles A (MxK)
* and B (KxN) can be added, where:
* M = ThreadItemsY
* N = ThreadItemsX
* K = sizeof(dp_vector_t) / sizeof(value_t).
*
* In order to leverage architecture-specific "dot-product accumulate" ISA
* operations, K is dictated by the thread_accumulator class in the form of
* the member-type dp_vector_t, which defines a K-component vector of value_t.
* The multiplicand inputs A and B are provided as arrays of dp_vector_t having
* extents ThreadItemsY and ThreadItemsX, respectively. (In the single
* component "dp1" scenario where dp_vector_t == value_t and thus K == 1, the
* multiplication is simply the outer product of two vectors.)
*
* The accumulators are zero-initialized in a two-phase process (construction +
* initialization) that requires shared storage in the form of the member-type
* scratch_storage_t during construction. (A single scratch_storage_t instance
* can be uniformly referenced across all threads in the block during
* construction *if* the block is synchronized between construction and
* initialization.)
*
* NB: This generic class is not directly constructible. Architecture- and
* algorithm-specific template specializations will provide the API
* functionality prescribed here.
*/
template <
int ThreadItemsY, ///< Height of thread tile in accum_t
int ThreadItemsX, ///< Width of thread tile in accum_t
typename value_t, ///< Multiplicand value type
typename accum_t, ///< Accumulator value type
int ACCUM_BYTES = ///< Size in bytes of accum_t
sizeof(accum_t),
arch_family_t::kind_t ArchFamily = ///< Architectural family enumerant
CUTLASS_ARCH_FAMILY>
struct thread_accumulator
{
protected:
//-------------------------------------------------------------------------
// Constants and types
//-------------------------------------------------------------------------
/// Specialized dot-product traits type
typedef dp_accummulate<value_t, accum_t> dp_accum_traits_t;
public:
//-------------------------------------------------------------------------
// Member types
//-------------------------------------------------------------------------
/// Dot-product vector type
typedef typename dp_accum_traits_t::dp_vector_t dp_vector_t;
/// Scratch storage layout
struct scratch_storage_t {};
protected:
//-------------------------------------------------------------------------
// Data members
//-------------------------------------------------------------------------
/// Thread's tile of accumulators
accum_t accumulators[ThreadItemsY][ThreadItemsX];
//-------------------------------------------------------------------------
// Utility methods
//-------------------------------------------------------------------------
/**
* Compute a multiply-add at accumulator coordinates (x, y)
*/
inline __device__
void mad_xy(
dp_vector_t (&tile_a)[ThreadItemsY],
dp_vector_t (&tile_b)[ThreadItemsX],
int x,
int y)
{
dp_accum_traits_t::mad(
accumulators[y][x],
tile_a[y],
tile_b[x],
accumulators[y][x]);
}
public:
//-------------------------------------------------------------------------
// Constructor API
//-------------------------------------------------------------------------
/// Constructor
inline __device__
thread_accumulator(
scratch_storage_t &scratch)
{}
//-------------------------------------------------------------------------
// Accumulator API
//-------------------------------------------------------------------------
/**
* \brief Zero-initialize thread accumulators.
*
* If a common reference to a single block-wide shared instance of scratch_storage_t
* is used during construction, the block must be synchronized after construction
* but prior to the invocation of init().
*/
inline __device__
void init()
{
#pragma unroll
for (int y = 0; y < ThreadItemsY; ++y) {
#pragma unroll
for (int x = 0; x < ThreadItemsX; ++x)
{
accumulators[y][x] = accum_t(0);
}
}
}
/**
* Retrieve the accumulator at thread tile coordinates (x, y)
*/
inline __device__
accum_t get(int x, int y)
{
// Accumulators are row-major
return accumulators[y][x];
}
/**
* \brief Compute the product of tile_a and tile_b and add the result to
* the tile of accumulators.
*/
inline __device__
void multiply_accumulate(
dp_vector_t (&tile_a)[ThreadItemsY],
dp_vector_t (&tile_b)[ThreadItemsX])
{
// Simply traverse the accumulator tile in row-major order
#pragma unroll
for (int y = 0; y < ThreadItemsY; ++y)
{
#pragma unroll
for (int x = 0; x < ThreadItemsX; ++x)
{
mad_xy(tile_a, tile_b, x, y);
}
}
}
};
/******************************************************************************
* thread_accumulator (__half->__half specialization)
******************************************************************************/
/**
* \brief Thread-level multiply-accumulate abstraction (__half->__half specialization)
*
* NB: Because we use the 2-item SIMD instruction HFMA2:
* - ThreadItemsX must be an even multiple of 2
* - ThreadItemsY must be an even multiple of 2
*
*/
template <
int ThreadItemsY, ///< Height in rows of thread tile in C
int ThreadItemsX, ///< Width in columns of thread tile in C
arch_family_t::kind_t ArchFamily> ///< Architectural family enumerant
struct thread_accumulator<
ThreadItemsY,
ThreadItemsX,
__half, ///< Multiplicand value type (matrices A and B)
__half, ///< Accumulator value type (matrix C and scalars)
2, ///< Size in bytes of accum_t
ArchFamily>
{
protected:
//-------------------------------------------------------------------------
// Constants and types
//-------------------------------------------------------------------------
/// Constants
enum
{
/// Height of thread tile in column-major uint32_t SIMD pairs along Y dimension
ThreadTilePairsY = divide_assert<ThreadItemsY, 2>::value,
/// Width of thread tile in column-major uint32_t SIMD pairs along X dimension
ThreadTilePairsX = ThreadItemsX,
/// Number of SIMD pairs in thread's slice of block-wide tile multiplicand A
ThreadPairsA = divide_assert<ThreadItemsY, 2>::value,
/// Number of SIMD pairs in thread's slice of block-wide tile multiplicand B
ThreadPairsB = divide_assert<ThreadItemsX, 2>::value,
};
public:
//-------------------------------------------------------------------------
// Member types
//-------------------------------------------------------------------------
/// Dot-product vector type
typedef __half dp_vector_t;
/// Scratch storage layout
struct scratch_storage_t {};
private:
//-------------------------------------------------------------------------
// Members
//-------------------------------------------------------------------------
/// Thread's tile of C accumulator pairs (the uint32_t SIMD pairs are
/// column-major, the 2D tile layout is also column-major)
uint32_t accumulator_pairs[ThreadTilePairsX][ThreadTilePairsY];
//-------------------------------------------------------------------------
// Utility methods
//-------------------------------------------------------------------------
/**
* Compute an HFMA2 MAD
*/
inline __device__ void mad(
uint32_t &d,
const uint32_t &a,
const uint32_t &b,
const uint32_t &c)
{
asm volatile ("fma.rn.f16x2 %0, %1, %2, %3;\n"
: "=r"(d) : "r"(a), "r"(b), "r"(c));
}
/**
* Compute an HFMA2 MAD with replicated b.lo:
* d{hi} = a{hi} * b{lo} + c{hi};
* d{lo} = a{lo} * b{lo} + c{lo};
*/
inline __device__ void mad_replicate_low(
uint32_t &d,
const uint32_t &a,
const uint32_t &b,
const uint32_t &c)
{
// Replicate low halves of b
uint32_t replicate;
asm volatile (
"{"
" .reg .b16 b_low,b_high;\n"
" mov.b32 {b_low,b_high}, %1;\n"
" mov.b32 %0, {b_low,b_low};\n"
"}" : "=r"(replicate) : "r"(b));
mad(d, a, replicate, c);
}
/**
* Compute an HFMA2 MAD with replicated b.hi:
* d{hi} = a{hi} * b{hi} + c{hi};
* d{lo} = a{lo} * b{hi} + c{lo};
*/
inline __device__ void mad_replicate_high(
uint32_t &d,
const uint32_t &a,
const uint32_t &b,
const uint32_t &c)
{
// Replicate high halves of b
uint32_t replicate;
asm volatile (
"{"
" .reg .b16 b_low,b_high;\n"
" mov.b32 {b_low,b_high}, %1;\n"
" mov.b32 %0, {b_high,b_high};\n"
"}" : "=r"(replicate) : "r"(b));
mad(d, a, replicate, c);
}
/**
* Compute a multiply-add at accumulator SIMD-pair coordinates (pair_x, pair_y)
*/
inline __device__
void mad_xy_even(
uint32_t (&pairs_tile_a)[ThreadPairsA],
uint32_t (&pairs_tile_b)[ThreadPairsB],
int pair_x,
int pair_y)
{
// Even column: use low half of the b pair
mad_replicate_low(
accumulator_pairs[pair_x][pair_y],
pairs_tile_a[pair_y],
pairs_tile_b[pair_x / 2],
accumulator_pairs[pair_x][pair_y]);
}
/**
* Compute a multiply-add at accumulator SIMD-pair coordinates (pair_x, pair_y)
*/
inline __device__
void mad_xy_odd(
uint32_t (&pairs_tile_a)[ThreadPairsA],
uint32_t (&pairs_tile_b)[ThreadPairsB],
int pair_x,
int pair_y)
{
// Odd column: use high half of the b pair
mad_replicate_high(
accumulator_pairs[pair_x][pair_y],
pairs_tile_a[pair_y],
pairs_tile_b[pair_x / 2],
accumulator_pairs[pair_x][pair_y]);
}
public:
//-------------------------------------------------------------------------
// Constructor API
//-------------------------------------------------------------------------
/// Constructor
inline __device__
thread_accumulator(
scratch_storage_t &scratch)
{}
//-------------------------------------------------------------------------
// Accumulator API
//-------------------------------------------------------------------------
/**
* Zero-initialize thread accumulators.
*/
inline __device__
void init()
{
#pragma unroll
for (int y = 0; y < ThreadTilePairsY; ++y)
{
#pragma unroll
for (int x = 0; x < ThreadTilePairsX; ++x)
{
accumulator_pairs[x][y] = 0;
}
}
}
/**
* Retrieve the accumulator at thread tile coordinates (x, y)
*/
inline __device__
__half get(int x, int y)
{
// SIMD pairs are column-major
uint32_t pair = accumulator_pairs[x][y / 2];
return reinterpret_cast<__half (&)[2]>(pair)[y % 2];
}
/**
* \brief Compute the product of pairs_tile_a and pairs_tile_b and add the result to
* the tile of accumulators.
*/
inline __device__
void multiply_accumulate(
dp_vector_t (&tile_a)[ThreadItemsY],
dp_vector_t (&tile_b)[ThreadItemsX])
{
typedef uint32_t pairs_tile_a_t[ThreadPairsA];
typedef uint32_t pairs_tile_b_t[ThreadPairsB];
// Alias slices in pairs
pairs_tile_a_t &pairs_tile_a = reinterpret_cast<pairs_tile_a_t&>(tile_a);
pairs_tile_b_t &pairs_tile_b = reinterpret_cast<pairs_tile_b_t&>(tile_b);
// Simply traverse the accumulator tile in column-major order
#pragma unroll
for (int x = 0; x < ThreadTilePairsX; ++x)
{
#pragma unroll
for (int y = 0; y < ThreadTilePairsY; ++y)
{
mad_xy_even(pairs_tile_a, pairs_tile_b, x, y);
}
}
}
};
} // namespace gemm
} // namespace cutlass

View File

@ -1,207 +0,0 @@
/******************************************************************************
* Copyright (c) 2011-2017, NVIDIA CORPORATION. All rights reserved.
*
* Redistribution and use in source and binary forms, with or without
* modification, are not permitted.
*
* THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
* ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
* WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
* DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE FOR ANY
* DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
* (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
* LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
* ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
* (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
* SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
*
******************************************************************************/
#pragma once
/**
* \file
* Thread-level multiply-accumulate abstraction
* (Volta 4B accum_t specialization)
*/
#include <mma.h>
#include "../util/util.h"
#include "dp_accummulate.h"
namespace cutlass {
namespace gemm {
/*!
*\brief matrix_layout to perform conversion between Cutlass types and WMMA types
*/
template <matrix_transform_t::kind_t>
struct matrix_layout;
/// Maps matrix_transform_t::NonTranspose to nvcuda::wmma::mem_col_major
template <>
struct matrix_layout<matrix_transform_t::NonTranspose>
{
/// Type tag in nvcuda::wmma namespace
typedef nvcuda::wmma::col_major tag;
/// Column major layout
static const nvcuda::wmma::layout_t kind = nvcuda::wmma::mem_col_major;
/// Cutlass matrix transform kind
static const matrix_transform_t::kind_t cutlass_kind = matrix_transform_t::NonTranspose;
};
/// Maps matrix_transform_t::NonTranspose to nvcuda::wmma::mem_row_major
template <>
struct matrix_layout<matrix_transform_t::Transpose>
{
/// Type tag in nvcuda::wmma namespace
typedef nvcuda::wmma::row_major tag;
/// Column major layout
static const nvcuda::wmma::layout_t kind = nvcuda::wmma::mem_row_major;
/// Cutlass matrix transform kind
static const matrix_transform_t::kind_t cutlass_kind = matrix_transform_t::Transpose;
};
/*!
* \brief Warp-synchronous matrix multiply-accumulate abstraction
*
* wmma_accumulator maps the CUDA WMMA API onto the GEMM structure
*/
template <
int WarpItemsY, /// Number of rows of the warp's accumulator tile
int WarpItemsX, /// Number of columns of the warp's accumulator tile
int WmmaItemsY, /// Number of rows in a single WMMA operation
int WmmaItemsX, /// Number of columns in a single WMMA operation
int WmmaItemsK, /// Inner dimension of WMMA operation
typename value_a_t, /// Type of A operand
typename value_b_t, /// Type of B operand
typename accum_t, /// Type of source and destination accumulators
matrix_transform_t::kind_t TransformA, /// Layout of A operand
matrix_transform_t::kind_t TransformB /// Layout of B operand
>
struct wmma_accumulator
{
public:
//-------------------------------------------------------------------------
// Constants and types
//-------------------------------------------------------------------------
enum
{
/// Number of WMMA blocks in warp row
WmmaBlocksX = divide_assert<WarpItemsX, WmmaItemsX>::value,
/// Number of WMMA blocks in a warp column
WmmaBlocksY = divide_assert<WarpItemsY, WmmaItemsY>::value,
};
/// Fragment type for matrix operand A
typedef nvcuda::wmma::fragment<
nvcuda::wmma::matrix_a,
WmmaItemsY,
WmmaItemsX,
WmmaItemsK,
value_a_t,
typename matrix_layout<TransformA>::tag>
fragment_a_t;
/// Fragment type for matrix operand B
typedef nvcuda::wmma::fragment<
nvcuda::wmma::matrix_b,
WmmaItemsY,
WmmaItemsX,
WmmaItemsK,
value_b_t,
typename matrix_layout<TransformB>::tag>
fragment_b_t;
/// Fragment type for accumulator
typedef nvcuda::wmma::fragment<
nvcuda::wmma::accumulator,
WmmaItemsY,
WmmaItemsX,
WmmaItemsK,
accum_t>
accumulator_t;
/// Scratch storage layout
struct scratch_storage_t
{
/// Initialization vector
uint4 zero_slab;
};
public:
//-------------------------------------------------------------------------
// Data members
//-------------------------------------------------------------------------
/// Thread's tile of accumulators
accumulator_t accumulators[WmmaBlocksX][WmmaBlocksY];
public:
//-------------------------------------------------------------------------
// Constructor API
//-------------------------------------------------------------------------
/// Constructor initializes accumulators to zero
inline __device__
wmma_accumulator()
{
init();
}
//-------------------------------------------------------------------------
// Accumulator API
//-------------------------------------------------------------------------
/**
* \brief Zero-initialize thread accumulators.
*/
inline __device__
void init()
{
#pragma unroll
for (int x = 0; x < WmmaBlocksX; ++x)
{
#pragma unroll
for (int y = 0; y < WmmaBlocksY; ++y)
{
nvcuda::wmma::fill_fragment(accumulators[x][y], accum_t(0));
}
}
}
/**
* \brief Compute the product of tile_a and tile_b and add the result to
* the tile of accumulators.
*/
inline __device__
void multiply_accumulate(
fragment_a_t (&tile_a)[WmmaBlocksY],
fragment_b_t (&tile_b)[WmmaBlocksX])
{
#pragma unroll
for (int x = 0; x < WmmaBlocksX; ++x)
{
#pragma unroll
for (int y = 0; y < WmmaBlocksY; ++y)
{
nvcuda::wmma::mma_sync(accumulators[x][y], tile_a[y], tile_b[x], accumulators[x][y]);
}
}
}
};
} // namespace gemm
} // namespace cutlass

View File

@ -1,112 +0,0 @@
/******************************************************************************
* Copyright (c) 2011-2017, NVIDIA CORPORATION. All rights reserved.
*
* Redistribution and use in source and binary forms, with or without
* modification, are not permitted.
*
* THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
* ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
* WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
* DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE FOR ANY
* DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
* (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
* LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
* ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
* (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
* SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
*
******************************************************************************/
#pragma once
/**
* \file
* \brief Debugging and logging functionality
*/
#include <stdio.h>
namespace cutlass {
/******************************************************************************
* Debug and logging macros
******************************************************************************/
/**
* Formats and prints the given message to stdout
*/
#if !defined(CUDA_LOG)
#if !defined(__CUDA_ARCH__)
#define CUDA_LOG(format, ...) printf(format,__VA_ARGS__)
#else
#define CUDA_LOG(format, ...) printf("[block (%d,%d,%d), thread (%d,%d,%d)]: " format, blockIdx.x, blockIdx.y, blockIdx.z, threadIdx.x, threadIdx.y, threadIdx.z, __VA_ARGS__);
#endif
#endif
/**
* Formats and prints the given message to stdout only if DEBUG is defined
*/
#if !defined(CUDA_LOG_DEBUG)
#ifdef DEBUG
#define CUDA_LOG_DEBUG(format, ...) CUDA_LOG(format, __VA_ARGS__)
#else
#define CUDA_LOG_DEBUG(format, ...)
#endif
#endif
/**
* \brief The corresponding error message is printed to \p stderr (or \p stdout in device code) along with the supplied source context.
*
* \return The CUDA error.
*/
__host__ __device__ inline cudaError_t cuda_perror_impl(
cudaError_t error,
const char* filename,
int line)
{
(void)filename;
(void)line;
if (error)
{
#if !defined(__CUDA_ARCH__)
fprintf(stderr, "CUDA error %d [%s, %d]: %s\n", error, filename, line, cudaGetErrorString(error));
fflush(stderr);
#else
printf("CUDA error %d [%s, %d]\n", error, filename, line);
#endif
}
return error;
}
/**
* \brief Perror macro
*/
#ifndef CUDA_PERROR
#define CUDA_PERROR(e) cuda_perror_impl((cudaError_t) (e), __FILE__, __LINE__)
#endif
/**
* \brief Perror macro with exit
*/
#ifndef CUDA_PERROR_EXIT
#define CUDA_PERROR_EXIT(e) if (cuda_perror_impl((cudaError_t) (e), __FILE__, __LINE__)) { exit(1); }
#endif
/**
* \brief Perror macro only if DEBUG is defined
*/
#ifndef CUDA_PERROR_DEBUG
#ifdef DEBUG
#define CUDA_PERROR_DEBUG(e) CUDA_PERROR(e)
#else
#define CUDA_PERROR_DEBUG(e) (e)
#endif
#endif
} // namespace cutlass

View File

@ -1,216 +0,0 @@
/******************************************************************************
* Copyright (c) 2011-2017, NVIDIA CORPORATION. All rights reserved.
*
* Redistribution and use in source and binary forms, with or without
* modification, are not permitted.
*
* THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
* ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
* WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
* DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE FOR ANY
* DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
* (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
* LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
* ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
* (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
* SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
*
******************************************************************************/
#pragma once
/**
* \file
* \brief Utilities for device introspection
*/
#include "debug.h"
#include "nv_std.h"
#include "printable.h"
namespace cutlass {
/******************************************************************************
* math_operation_class_t
*
* Enumeration to select the appropriate math operation
*
* The assumption is multiple math operations may be used to compute GEMM
* for a given selection of operand and accumulator types.
*
******************************************************************************/
/// Math operation
enum class math_operation_class_t
{
scalar, // scalar (and vector) multiply-accumulate operations
matrix // Volta tensor operations
};
/******************************************************************************
* arch_family_t
******************************************************************************/
/**
* \brief Enumeration of NVIDIA GPU architectural families
*/
struct arch_family_t
{
/// \brief Enumerants
enum kind_t
{
Unsupported = 0,
Kepler = 3,
Maxwell = 5,
Volta = 7,
};
/// Enumerant value
kind_t kind;
/// Default constructor
arch_family_t() : kind(Unsupported) {}
/// Copy constructor
arch_family_t(const kind_t &other_kind) : kind(other_kind) {}
/// Cast to kind_t
operator kind_t() const { return kind; }
/// Returns the instance as a string
__host__ __device__ inline
char const* to_string() const
{
switch (kind)
{
case Kepler: return "Kepler";
case Maxwell: return "Maxwell";
case Volta: return "Volta";
case Unsupported:
default: return "Unsupported";
}
}
/// Insert the formatted instance into the output stream
void print(std::ostream& out) const { out << to_string(); }
};
/**
* Macro for architecture targeted by the current compiler pass
*/
#if defined(__CUDA_ARCH__)
#define CUTLASS_ARCH __CUDA_ARCH__
#else
#define CUTLASS_ARCH 0
#endif
/**
* Macro for architecture family targeted by the current compiler pass
*/
#define CUTLASS_ARCH_FAMILY \
( \
(CUTLASS_ARCH < 300) ? \
arch_family_t::Unsupported : \
(CUTLASS_ARCH < 500) ? \
arch_family_t::Kepler : \
(CUTLASS_ARCH < 700) ? \
arch_family_t::Maxwell : \
arch_family_t::Volta \
)
/******************************************************************************
* Device introspection
******************************************************************************/
/**
* Empty kernel for querying PTX manifest metadata (e.g., version) for the current device
*/
template <typename T>
__global__ void empty_kernel(void) { }
/**
* \brief Retrieves the PTX version that will be used on the current device (major * 100 + minor * 10)
*/
cudaError_t ptx_version(int &version)
{
struct Dummy
{
/// Type definition of the empty_kernel kernel entry point
typedef void (*EmptyKernelPtr)();
/// Force empty_kernel<void> to be generated if this class is used
EmptyKernelPtr Empty()
{
return empty_kernel<void>;
}
};
cudaError_t error = cudaSuccess;
do
{
cudaFuncAttributes empty_kernel_attrs;
if (CUDA_PERROR_DEBUG(error = cudaFuncGetAttributes(&empty_kernel_attrs, empty_kernel<void>))) break;
version = empty_kernel_attrs.ptxVersion * 10;
}
while (0);
return error;
}
/**
* \brief Retrieves the SM version (major * 100 + minor * 10) for the current device
*/
cudaError_t get_sm_version(int &sm_version)
{
cudaError_t error = cudaSuccess;
// Get device ordinal
int device_ordinal;
if (CUDA_PERROR_DEBUG(error = cudaGetDevice(&device_ordinal)))
return error;
// Fill in SM version
int major, minor;
if (CUDA_PERROR_DEBUG(error = cudaDeviceGetAttribute(&major, cudaDevAttrComputeCapabilityMajor, device_ordinal)))
return error;
if (CUDA_PERROR_DEBUG(error = cudaDeviceGetAttribute(&minor, cudaDevAttrComputeCapabilityMinor, device_ordinal)))
return error;
sm_version = major * 100 + minor * 10;
return error;
}
/**
* \brief Retrieves the count for the current device
*/
cudaError_t get_sm_count(int &sm_count)
{
cudaError_t error = cudaSuccess;
// Get device ordinal
int device_ordinal;
if (CUDA_PERROR_DEBUG(error = cudaGetDevice(&device_ordinal)))
return error;
// Get SM count
if (CUDA_PERROR_DEBUG(error = cudaDeviceGetAttribute (&sm_count, cudaDevAttrMultiProcessorCount, device_ordinal)))
return error;
return error;
}
} // namespace cutlass

View File

@ -1,484 +0,0 @@
/******************************************************************************
* Copyright (c) 2011-2017, NVIDIA CORPORATION. All rights reserved.
*
* Redistribution and use in source and binary forms, with or without
* modification, are not permitted.
*
* THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
* ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
* WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
* DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE FOR ANY
* DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
* (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
* LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
* ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
* (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
* SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
*
******************************************************************************/
#pragma once
/**
* \file
* \brief I/O device intrinsics
*/
#include <stdint.h>
#include <cuda_fp16.h>
#include "nv_std.h"
#include "math.h"
namespace cutlass {
/******************************************************************************
* io_vector
******************************************************************************/
/**
* Base aligned storage for IO vector
*/
template <typename value_t, int VectorItems, int AlignBytes> struct io_vector_base;
template <typename value_t, int VectorItems> struct __align__(1) io_vector_base<value_t, VectorItems, 1> { value_t buff[VectorItems]; };
template <typename value_t, int VectorItems> struct __align__(2) io_vector_base<value_t, VectorItems, 2> { value_t buff[VectorItems]; };
template <typename value_t, int VectorItems> struct __align__(4) io_vector_base<value_t, VectorItems, 4> { value_t buff[VectorItems]; };
template <typename value_t, int VectorItems> struct __align__(8) io_vector_base<value_t, VectorItems, 8> { value_t buff[VectorItems]; };
template <typename value_t, int VectorItems> struct __align__(16) io_vector_base<value_t, VectorItems, 16> { value_t buff[VectorItems]; };
/**
* \brief Aligned vector type for coarsening data movement instructions
*
* Exposes the member constant \p VectorItems, the actual number of component
* values comprising the io_vector
*/
template <
typename value_t, ///< Component value type
int MaxVectorItems, ///< Maximum allowable component values
int MaxAlignBytes ///< Maximum allowable alignment
= __NV_STD_MIN(16, MaxVectorItems * sizeof(value_t)),
int AlignBytes ///< Actual alignment
= __NV_STD_MIN(sizeof(value_t) * MaxVectorItems, MaxAlignBytes),
int VectorItems ///< Actual number of component values
= divide_assert<AlignBytes, sizeof(value_t)>::value,
bool MustAlias ///< Whether we need to alias during loads/stores
= (VectorItems > 4)>
struct io_vector;
/**
* IO vector (specialization for VectorItems <= 4)
*/
template <
typename value_t,
int MaxVectorItems,
int MaxAlignBytes,
int _AlignBytes,
int _VectorItems>
struct io_vector <
value_t,
MaxVectorItems,
MaxAlignBytes,
_AlignBytes,
_VectorItems,
false>
:
io_vector_base<value_t, _VectorItems, _AlignBytes>
{
enum
{
VectorItems = _VectorItems,
AlignBytes = _AlignBytes
};
static_assert(is_pow2<AlignBytes>::value, "I/O vector alignment must be a power-of-two.");
static_assert((AlignBytes <= 16), "I/O vector alignment must <= 16B.");
inline __device__
void load(const io_vector *ptr)
{
*this = *ptr;
}
inline __device__
void load(const value_t *ptr)
{
*this = *reinterpret_cast<const io_vector*>(ptr);
}
inline __device__
void store(io_vector *ptr) const
{
*ptr = *this;
}
inline __device__
void store(value_t *ptr) const
{
*reinterpret_cast<io_vector*>(ptr) = *this;
}
};
/**
* IO vector (specialization for VectorItems > 4)
*
* NB: Workaround for NVCC not generating 128-bit loads/stores for aligned
* structures having component types < 32b
*/
template <
typename value_t,
int MaxVectorItems,
int MaxAlignBytes,
int _AlignBytes,
int _VectorItems>
struct io_vector <
value_t,
MaxVectorItems,
MaxAlignBytes,
_AlignBytes,
_VectorItems,
true>
:
io_vector_base<value_t, _VectorItems, _AlignBytes>
{
enum
{
VectorItems = _VectorItems,
AlignBytes = _AlignBytes
};
static_assert(is_pow2<AlignBytes>::value, "I/O vector alignment must be a power-of-two.");
static_assert((AlignBytes <= 16), "I/O vector alignment must <= 16B.");
typedef typename nv_std::conditional<(AlignBytes == 8),
uint2, // Use 8B load
uint4> // Use 16B load
::type align_t;
inline __device__
void load(const io_vector *ptr)
{
*reinterpret_cast<align_t*>(this) = *reinterpret_cast<const align_t*>(ptr);
}
inline __device__
void load(const value_t *ptr)
{
*reinterpret_cast<align_t*>(this) = *reinterpret_cast<const align_t*>(ptr);
}
inline __device__
void store(io_vector *ptr) const
{
*reinterpret_cast<align_t*>(ptr) = *reinterpret_cast<const align_t*>(this);
}
inline __device__
void store(value_t *ptr) const
{
*reinterpret_cast<align_t*>(ptr) = *reinterpret_cast<const align_t*>(this);
}
};
/******************************************************************************
* Macro expansions for vector loads
******************************************************************************/
/**
* Define vector-4 LD specialization for the given load modifier
*/
#define CUTLASS_LD_V4(f_name, value_t, load_modifier, ptx_type, val_constraint, ptr_constraint) \
template <typename ptr_t> \
inline __device__ \
void f_name( \
value_t (&dest)[4], \
ptr_t ptr) \
{ \
asm volatile ("ld."#load_modifier".v4."#ptx_type" {%0, %1, %2, %3}, [%4];\n" \
: \
"="#val_constraint(dest[0]), \
"="#val_constraint(dest[1]), \
"="#val_constraint(dest[2]), \
"="#val_constraint(dest[3]) \
: \
#ptr_constraint(ptr)); \
}
/**
* Define vector-2 LD specialization for the given load modifier
*/
#define CUTLASS_LD_V2(f_name, value_t, load_modifier, ptx_type, val_constraint, ptr_constraint) \
template <typename ptr_t> \
inline __device__ \
void f_name( \
value_t (&dest)[2], \
ptr_t ptr) \
{ \
asm volatile ("ld."#load_modifier".v2."#ptx_type" {%0, %1}, [%2];\n" \
: \
"="#val_constraint(dest[0]), \
"="#val_constraint(dest[1]) \
: \
#ptr_constraint(ptr)); \
}
/**
* Define vector-1 LD specialization for the given load modifier
*/
#define CUTLASS_LD_V1(f_name, value_t, load_modifier, ptx_type, val_constraint, ptr_constraint) \
template <typename ptr_t> \
inline __device__ \
void f_name( \
value_t (&dest)[1], \
ptr_t ptr) \
{ \
asm volatile ("ld."#load_modifier"."#ptx_type" %0, [%1];\n" \
: \
"="#val_constraint(dest[0]) \
: \
#ptr_constraint(ptr)); \
}
/**
* Define powers-of-two vector LD specializations
*/
#define CUTLASS_LD_ALL(f_name, value_t, load_modifier, ptx_type, val_constraint, ptr_constraint) \
CUTLASS_LD_V4(f_name, value_t, load_modifier, ptx_type, val_constraint, ptr_constraint) \
CUTLASS_LD_V2(f_name, value_t, load_modifier, ptx_type, val_constraint, ptr_constraint) \
CUTLASS_LD_V1(f_name, value_t, load_modifier, ptx_type, val_constraint, ptr_constraint)
/******************************************************************************
* Macro expansions for vector stores
******************************************************************************/
/**
* Define vector-4 ST specialization for the given load modifier
*/
#define CUTLASS_ST_V4(f_name, value_t, store_modifier, ptx_type, val_constraint, ptr_constraint) \
template <typename ptr_t> \
inline __device__ \
void f_name( \
ptr_t ptr, \
const value_t (&src)[4]) \
{ \
asm volatile ("st."#store_modifier".v4."#ptx_type" [%0], {%1, %2, %3, %4};\n" \
: : \
#ptr_constraint(ptr), \
#val_constraint(src[0]), \
#val_constraint(src[1]), \
#val_constraint(src[2]), \
#val_constraint(src[3])); \
}
/**
* Define vector-2 ST specialization for the given load modifier
*/
#define CUTLASS_ST_V2(f_name, value_t, store_modifier, ptx_type, val_constraint, ptr_constraint) \
template <typename ptr_t> \
inline __device__ \
void f_name( \
ptr_t ptr, \
const value_t (&src)[2]) \
{ \
asm volatile ("st."#store_modifier".v2."#ptx_type" [%0], {%1, %2};\n" \
: : \
#ptr_constraint(ptr), \
#val_constraint(src[0]), \
#val_constraint(src[1])); \
}
/**
* Define vector-1 ST specialization for the given load modifier
*/
#define CUTLASS_ST_V1(f_name, value_t, store_modifier, ptx_type, val_constraint, ptr_constraint) \
template <typename ptr_t> \
inline __device__ \
void f_name( \
ptr_t ptr, \
const value_t (&src)[1]) \
{ \
asm volatile ("st."#store_modifier"."#ptx_type" [%0], %1;\n" \
: : \
#ptr_constraint(ptr), \
#val_constraint(src[0])); \
}
/**
* Define powers-of-two vector LD specializations
*/
#define CUTLASS_ST_ALL(f_name, value_t, load_modifier, ptx_type, val_constraint, ptr_constraint) \
CUTLASS_ST_V4(f_name, value_t, load_modifier, ptx_type, val_constraint, ptr_constraint) \
CUTLASS_ST_V2(f_name, value_t, load_modifier, ptx_type, val_constraint, ptr_constraint) \
CUTLASS_ST_V1(f_name, value_t, load_modifier, ptx_type, val_constraint, ptr_constraint)
/******************************************************************************
* Macro expansions for vector IO
******************************************************************************/
/**
* Define global and shared LD specializations
*/
#define CUTLASS_IO(value_t, ptx_type, val_constraint) \
CUTLASS_LD_ALL(ldg_cg_internal, value_t, global.cg, ptx_type, val_constraint, l) \
CUTLASS_ST_ALL(stg_cg_internal, value_t, global.cg, ptx_type, val_constraint, l)
// Define IO for useful types
CUTLASS_IO(double, f64, d)
CUTLASS_IO(float, f32, f)
CUTLASS_IO(int64_t, b64, l)
CUTLASS_IO(int32_t, b32, r)
CUTLASS_IO(int16_t, b16, h)
// Macro cleanup
#undef CUTLASS_IO
#undef CUTLASS_LD_ALL
#undef CUTLASS_LD_V4
#undef CUTLASS_LD_V2
#undef CUTLASS_LD_V1
#undef CUTLASS_ST_ALL
#undef CUTLASS_ST_V4
#undef CUTLASS_ST_V2
#undef CUTLASS_ST_V1
/******************************************************************************
* I/O cast types
******************************************************************************/
/// Provides the type for which to reinterpret-cast a given vector
template <
typename value_t,
int IoVecDim,
int ValueBytes = sizeof(value_t)>
struct io_cast
{
typedef value_t type[IoVecDim];
};
/// Provides the type for which to reinterpret-cast a vector of 1B types
template <
typename value_t,
int IoVecDim>
struct io_cast<value_t, IoVecDim, 1>
{
typedef typename nv_std::conditional<
(IoVecDim < 2),
int8_t[1], // Use 8b load
typename nv_std::conditional<
(IoVecDim < 4),
int16_t[1], // Use 16b load
int32_t[IoVecDim / 4]>::type>::type // Use up to 128b load
type;
};
/// Provides the type for which to reinterpret-cast a vector of 2B types
template <
typename value_t,
int IoVecDim>
struct io_cast<value_t, IoVecDim, 2>
{
typedef typename nv_std::conditional<
(IoVecDim < 2),
int16_t[1], // Use 16b load
int32_t[IoVecDim / 2]>::type // Use up to 128b load
type;
};
/******************************************************************************
* ldg_cg intrinsics
******************************************************************************/
/// Load from global (cache-global modifier)
template <typename value_t, typename ptr_t>
inline __device__
void ldg_cg(
value_t &dest,
ptr_t d_in)
{
// Cast dest to a different array type if necessary
ldg_cg_internal(
reinterpret_cast<typename io_cast<value_t, 1>::type &>(dest),
d_in);
}
/// Load from global (cache-global modifier)
template <typename value_t, int IoVecDim, typename ptr_t>
inline __device__
void ldg_cg(
value_t (&dest)[IoVecDim],
ptr_t d_in)
{
static_assert(is_pow2<IoVecDim>::value, "I/O vectors must be a power-of-two.");
// Cast dest to a different array type if necessary
ldg_cg_internal(
reinterpret_cast<typename io_cast<value_t, IoVecDim>::type &>(dest),
d_in);
}
/******************************************************************************
* stg_cg intrinsics
******************************************************************************/
/// Store to global (cache-global modifier)
template <typename ptr_t, typename value_t>
inline __device__
void stg_cg(
ptr_t dest,
const value_t &src)
{
// Cast src to a different array type if necessary
stg_cg_internal(
dest,
reinterpret_cast<const typename io_cast<value_t, 1>::type &>(src));
}
/// Store to global (cache-global modifier)
template <typename ptr_t, int IoVecDim, typename value_t>
inline __device__
void stg_cg(
ptr_t dest,
const value_t (&src)[IoVecDim])
{
static_assert(is_pow2<IoVecDim>::value, "I/O vectors must be a power-of-two.");
// Cast src to a different array type if necessary
stg_cg_internal(
dest,
reinterpret_cast<const typename io_cast<value_t, IoVecDim>::type &>(src));
}
} // namespace cutlass

View File

@ -1,189 +0,0 @@
/*
* Copyright 1993-2017 NVIDIA Corporation. All rights reserved.
*
* NOTICE TO LICENSEE:
*
* This source code and/or documentation ("Licensed Deliverables") are
* subject to NVIDIA intellectual property rights under U.S. and
* international Copyright laws.
*
* These Licensed Deliverables contained herein is PROPRIETARY and
* CONFIDENTIAL to NVIDIA and is being provided under the terms and
* conditions of a form of NVIDIA software license agreement by and
* between NVIDIA and Licensee ("License Agreement") or electronically
* accepted by Licensee. Notwithstanding any terms or conditions to
* the contrary in the License Agreement, reproduction or disclosure
* of the Licensed Deliverables to any third party without the express
* written consent of NVIDIA is prohibited.
*
* NOTWITHSTANDING ANY TERMS OR CONDITIONS TO THE CONTRARY IN THE
* LICENSE AGREEMENT, NVIDIA MAKES NO REPRESENTATION ABOUT THE
* SUITABILITY OF THESE LICENSED DELIVERABLES FOR ANY PURPOSE. IT IS
* PROVIDED "AS IS" WITHOUT EXPRESS OR IMPLIED WARRANTY OF ANY KIND.
* NVIDIA DISCLAIMS ALL WARRANTIES WITH REGARD TO THESE LICENSED
* DELIVERABLES, INCLUDING ALL IMPLIED WARRANTIES OF MERCHANTABILITY,
* NONINFRINGEMENT, AND FITNESS FOR A PARTICULAR PURPOSE.
* NOTWITHSTANDING ANY TERMS OR CONDITIONS TO THE CONTRARY IN THE
* LICENSE AGREEMENT, IN NO EVENT SHALL NVIDIA BE LIABLE FOR ANY
* SPECIAL, INDIRECT, INCIDENTAL, OR CONSEQUENTIAL DAMAGES, OR ANY
* DAMAGES WHATSOEVER RESULTING FROM LOSS OF USE, DATA OR PROFITS,
* WHETHER IN AN ACTION OF CONTRACT, NEGLIGENCE OR OTHER TORTIOUS
* ACTION, ARISING OUT OF OR IN CONNECTION WITH THE USE OR PERFORMANCE
* OF THESE LICENSED DELIVERABLES.
*
* U.S. Government End Users. These Licensed Deliverables are a
* "commercial item" as that term is defined at 48 C.F.R. 2.101 (OCT
* 1995), consisting of "commercial computer software" and "commercial
* computer software documentation" as such terms are used in 48
* C.F.R. 12.212 (SEPT 1995) and is provided to the U.S. Government
* only as a commercial end item. Consistent with 48 C.F.R.12.212 and
* 48 C.F.R. 227.7202-1 through 227.7202-4 (JUNE 1995), all
* U.S. Government End Users acquire the Licensed Deliverables with
* only those rights set forth herein.
*
* Any use of the Licensed Deliverables in individual and commercial
* software must include, in the user documentation and internal
* comments to the code, the above Disclaimer and U.S. Government End
* Users Notice.
*/
#pragma once
/**
* \file
* \brief Math utilities
*/
#include "nv_std.h"
namespace cutlass {
/******************************************************************************
* Static math utilities
******************************************************************************/
/**
* Statically determine if N is a power-of-two
*/
template <int N>
struct is_pow2 : nv_std::integral_constant<bool, (N & (N - 1)) == 0>
{};
/**
* Statically determine log2(N), rounded down
*/
template <int N, int CurrentVal = N, int Count = 0>
struct log2_down
{
/// Static logarithm value
enum { value = log2_down<N, (CurrentVal >> 1), Count + 1>::value };
};
// Base case
template <int N, int Count>
struct log2_down<N, 1, Count>
{
enum { value = Count };
};
/**
* Statically determine log2(N), rounded up
*/
template <int N, int CurrentVal = N, int Count = 0>
struct log2_up
{
/// Static logarithm value
enum { value = log2_up<N, (CurrentVal >> 1), Count + 1>::value };
};
// Base case
template <int N, int Count>
struct log2_up<N, 1, Count>
{
enum { value = ((1 << Count) < N) ? Count + 1 : Count };
};
/**
* Statically estimate sqrt(N) to the nearest power-of-two
*/
template <int N>
struct sqrt_est
{
enum { value = 1 << (log2_up<N>::value / 2) };
};
/**
* For performing a constant-division with a compile-time assertion that the
* Divisor evenly-divides the Dividend.
*/
template <int Dividend, int Divisor>
struct divide_assert
{
enum { value = Dividend / Divisor};
static_assert((Dividend % Divisor == 0), "Not an even multiple");
};
/******************************************************************************
* Rounding
******************************************************************************/
/**
* Round dividend up to the nearest multiple of divisor
*/
template <typename dividend_t, typename divisor_t>
inline __host__ __device__
dividend_t round_nearest(dividend_t dividend, divisor_t divisor)
{
return ((dividend + divisor - 1) / divisor) * divisor;
}
/**
* Greatest common divisor
*/
template <typename value_t>
inline __host__ __device__
value_t gcd(value_t a, value_t b)
{
for (;;)
{
if (a == 0) return b;
b %= a;
if (b == 0) return a;
a %= b;
}
}
/**
* Least common multiple
*/
template <typename value_t>
inline __host__ __device__
value_t lcm(value_t a, value_t b)
{
value_t temp = gcd(a, b);
return temp ? (a / temp * b) : 0;
}
} // namespace cutlass

View File

@ -1,94 +0,0 @@
/******************************************************************************
* Copyright (c) 2011-2017, NVIDIA CORPORATION. All rights reserved.
*
* Redistribution and use in source and binary forms, with or without
* modification, are not permitted.
*
* THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
* ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
* WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
* DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE FOR ANY
* DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
* (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
* LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
* ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
* (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
* SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
*
******************************************************************************/
#pragma once
/**
* \file
* \brief Enumeration of dense matrix view transformations
*/
#include "printable.h"
namespace cutlass {
/******************************************************************************
* matrix_transform_t
******************************************************************************/
/**
* \brief Enumeration of dense matrix view transformations
*
* These enumerators (and corresponding tag types) describe which view
* transformation needs to be applied prior to operation upon a given dense
* matrix. Its values correspond to Fortran characters 'n' (non-transpose),
* 't'(transpose) and 'c'(conjugate transpose) that are often
* used as parameters to legacy BLAS implementations
*/
struct matrix_transform_t : printable_t
{
/// \brief Enumerants (same as CUBLAS)
enum kind_t
{
/// Invalid view
Invalid = -1,
/// Non-transpose view
NonTranspose = 0,
/// Transpose view
Transpose = 1,
/// Conjugate transpose view
ConjugateTranpose = 2,
};
/// Enumerant value
kind_t kind;
/// Default constructor
matrix_transform_t() : kind(Invalid) {}
/// Copy constructor
matrix_transform_t(const kind_t &other_kind) : kind(other_kind) {}
/// Cast to kind_t
operator kind_t() const { return kind; }
/// Returns the instance as a string
__host__ __device__ inline
char const* to_string() const
{
switch (kind)
{
case NonTranspose: return "NonTranspose";
case Transpose: return "Transpose";
case ConjugateTranpose: return "ConjugateTranpose";
default: return "Invalid";
}
}
/// Insert the formatted instance into the output stream
void print(std::ostream& out) const { out << to_string(); }
};
} // namespace cutlass

View File

@ -1,727 +0,0 @@
/*
* Copyright 1993-2017 NVIDIA Corporation. All rights reserved.
*
* NOTICE TO LICENSEE:
*
* This source code and/or documentation ("Licensed Deliverables") are
* subject to NVIDIA intellectual property rights under U.S. and
* international Copyright laws.
*
* These Licensed Deliverables contained herein is PROPRIETARY and
* CONFIDENTIAL to NVIDIA and is being provided under the terms and
* conditions of a form of NVIDIA software license agreement by and
* between NVIDIA and Licensee ("License Agreement") or electronically
* accepted by Licensee. Notwithstanding any terms or conditions to
* the contrary in the License Agreement, reproduction or disclosure
* of the Licensed Deliverables to any third party without the express
* written consent of NVIDIA is prohibited.
*
* NOTWITHSTANDING ANY TERMS OR CONDITIONS TO THE CONTRARY IN THE
* LICENSE AGREEMENT, NVIDIA MAKES NO REPRESENTATION ABOUT THE
* SUITABILITY OF THESE LICENSED DELIVERABLES FOR ANY PURPOSE. IT IS
* PROVIDED "AS IS" WITHOUT EXPRESS OR IMPLIED WARRANTY OF ANY KIND.
* NVIDIA DISCLAIMS ALL WARRANTIES WITH REGARD TO THESE LICENSED
* DELIVERABLES, INCLUDING ALL IMPLIED WARRANTIES OF MERCHANTABILITY,
* NONINFRINGEMENT, AND FITNESS FOR A PARTICULAR PURPOSE.
* NOTWITHSTANDING ANY TERMS OR CONDITIONS TO THE CONTRARY IN THE
* LICENSE AGREEMENT, IN NO EVENT SHALL NVIDIA BE LIABLE FOR ANY
* SPECIAL, INDIRECT, INCIDENTAL, OR CONSEQUENTIAL DAMAGES, OR ANY
* DAMAGES WHATSOEVER RESULTING FROM LOSS OF USE, DATA OR PROFITS,
* WHETHER IN AN ACTION OF CONTRACT, NEGLIGENCE OR OTHER TORTIOUS
* ACTION, ARISING OUT OF OR IN CONNECTION WITH THE USE OR PERFORMANCE
* OF THESE LICENSED DELIVERABLES.
*
* U.S. Government End Users. These Licensed Deliverables are a
* "commercial item" as that term is defined at 48 C.F.R. 2.101 (OCT
* 1995), consisting of "commercial computer software" and "commercial
* computer software documentation" as such terms are used in 48
* C.F.R. 12.212 (SEPT 1995) and is provided to the U.S. Government
* only as a commercial end item. Consistent with 48 C.F.R.12.212 and
* 48 C.F.R. 227.7202-1 through 227.7202-4 (JUNE 1995), all
* U.S. Government End Users acquire the Licensed Deliverables with
* only those rights set forth herein.
*
* Any use of the Licensed Deliverables in individual and commercial
* software must include, in the user documentation and internal
* comments to the code, the above Disclaimer and U.S. Government End
* Users Notice.
*/
#pragma once
/**
* \file
* \brief C++ features that may be otherwise unimplemented for CUDA device functions.
*
* This file has three components:
*
* (1) Macros:
* - Empty macro defines for C++ keywords not supported by the current
* version of C++. These simply allow compilation to proceed (but do
* not provide the added semantics).
* - \p noexcept
* - \p constexpr
* - \p nullptr
* - \p static_assert
*
* - Macro functions that we need in constant expressions because the
* C++ equivalents require constexpr compiler support. These are
* prefixed with \p __NV_STD_*
* - \p __NV_STD_MAX
* - \p __NV_STD_MIN
*
* (2) Re-implementations of STL functions and types:
* - C++ features that need the \p __device__ annotation. These are
* placed into the \p nv_std namespace.
* - \p plus
* - \p less
* - \p greater
* - \p min
* - \p max
* - \p methods on std::pair (==, !=, <, <=, >, >=, and make_pair())
*
* (3) Stop-gap implementations of unsupported STL functions and types:
* - STL functions and types defined by C++ 11/14/17/etc. that are not
* provided by the current version of C++. These are placed into the
* \p nv_std namespace
* - \p integral_constant
* - \p nullptr_t
* - \p true_type
* - \p false_type
* - \p bool_constant
* - \p enable_if
* - \p conditional
* - \p is_same
* - \p is_base_of
* - \p remove_const
* - \p remove_volatile
* - \p remove_cv
* - \p is_volatile
* - \p is_pointer
* - \p is_void
* - \p is_integral
* - \p is_floating_point
* - \p is_arithmetic
* - \p is_fundamental
* - \p is_trivially_copyable
* - \p alignment_of
* - \p aligned_storage
*
* (4) Functions and types that are STL-like (but aren't in the STL):
* - \p TODO: min and max functors?
*
* The idea is that, as we drop support for older compilers, we can simply #define
* the \p __NV_STD_XYZ macros and \p nv_std namespace to alias their C++
* counterparts (or trivially find-and-replace their occurrences in code text).
*/
//-----------------------------------------------------------------------------
// Include STL files that nv_std provides functionality for
//-----------------------------------------------------------------------------
#include <cstddef> // nullptr_t
#include <algorithm> // Minimum/maximum operations
#include <functional> // Arithmetic operations
#include <utility> // For methods on std::pair
#if (!defined(_MSC_VER) && (__cplusplus >= 201103L)) || (defined(_MSC_VER) && (_MS_VER >= 1500))
#include <type_traits> // For integral constants, conditional metaprogramming, and type traits
#endif
/******************************************************************************
* Macros
******************************************************************************/
//-----------------------------------------------------------------------------
// Keywords
//-----------------------------------------------------------------------------
/// noexcept, constexpr
#if (!defined(_MSC_VER) && (__cplusplus < 201103L)) || (defined(_MSC_VER) && (_MSC_VER < 1900))
#ifndef noexcept
#define noexcept
#endif
#ifndef constexpr
#define constexpr
#endif
#endif
/// nullptr
#if (!defined(_MSC_VER) && (__cplusplus < 201103L)) || (defined(_MSC_VER) && (_MSC_VER < 1310 ))
#ifndef nullptr
#define nullptr 0
#endif
#endif
/// static_assert
#if (!defined(_MSC_VER) && (__cplusplus < 201103L)) || (defined(_MSC_VER) && (_MSC_VER < 1600 ))
#ifndef static_assert
#define __nv_std_cat_(a, b) a ## b
#define __nv_std_cat(a, b) __nv_std_cat_(a, b)
#define static_assert(__e, __m) typedef int __nv_std_cat(AsSeRt, __LINE__)[(__e) ? 1 : -1]
#endif
#endif
//-----------------------------------------------------------------------------
// Functions
//-----------------------------------------------------------------------------
/// Select maximum(a, b)
#ifndef __NV_STD_MAX
#define __NV_STD_MAX(a, b) (((b) > (a)) ? (b) : (a))
#endif
/// Select minimum(a, b)
#ifndef __NV_STD_MIN
#define __NV_STD_MIN(a, b) (((b) < (a)) ? (b) : (a))
#endif
/******************************************************************************
* Re-implementations
******************************************************************************/
namespace nv_std {
//-----------------------------------------------------------------------------
// Arithmetic operations, comparisons <functional>
//-----------------------------------------------------------------------------
/// nv_std::plus
template <typename T>
struct plus
{
inline __host__ __device__
constexpr T operator()(const T &lhs, const T &rhs) const
{
return lhs + rhs;
}
};
/// std::less
template <typename T>
struct less
{
inline __host__ __device__
constexpr bool operator()(const T &lhs, const T &rhs) const
{
return lhs < rhs;
}
};
/// std::greater
template <typename T>
struct greater
{
inline __host__ __device__
constexpr bool operator()(const T &lhs, const T &rhs) const
{
return lhs > rhs;
}
};
//-----------------------------------------------------------------------------
// Minimum/maximum operations <algorithm>
//-----------------------------------------------------------------------------
/// std::min
template <typename T>
inline __host__ __device__
constexpr const T& min(
const T& a,
const T& b)
{
return (b < a) ? b : a;
}
/// std::max
template <typename T>
inline __host__ __device__
constexpr const T& max(
const T& a,
const T& b)
{
return (a < b) ? b : a;
}
//-----------------------------------------------------------------------------
// Methods on std::pair
//-----------------------------------------------------------------------------
using std::pair;
template< class T1, class T2 >
inline __host__ __device__
constexpr bool operator==( const pair<T1,T2>& lhs, const pair<T1,T2>& rhs )
{
return (lhs.first == rhs.first) && (lhs.second == rhs.second);
}
template< class T1, class T2 >
inline __host__ __device__
constexpr bool operator!=( const pair<T1,T2>& lhs, const pair<T1,T2>& rhs )
{
return (lhs.first != rhs.first) && (lhs.second != rhs.second);
}
template< class T1, class T2 >
inline __host__ __device__
constexpr bool operator<( const pair<T1,T2>& lhs, const pair<T1,T2>& rhs )
{
return (lhs.first < rhs.first) ?
true :
(rhs.first < lhs.first) ?
false :
(lhs.second < rhs.second);
}
template< class T1, class T2 >
inline __host__ __device__
constexpr bool operator<=( const pair<T1,T2>& lhs, const pair<T1,T2>& rhs )
{
return !(rhs < lhs);
}
template< class T1, class T2 >
inline __host__ __device__
constexpr bool operator>( const pair<T1,T2>& lhs, const pair<T1,T2>& rhs )
{
return (rhs < lhs);
}
template< class T1, class T2 >
inline __host__ __device__
constexpr bool operator>=( const pair<T1,T2>& lhs, const pair<T1,T2>& rhs )
{
return !(lhs < rhs);
}
template< class T1, class T2 >
inline __host__ __device__
std::pair<T1,T2> make_pair( T1 t, T2 u )
{
std::pair<T1,T2> retval;
retval.first = t;
retval.second = u;
return retval;
}
} // namespace nv_std
/******************************************************************************
* Implementations of C++ 11/14/17/... STL features
******************************************************************************/
namespace nv_std {
//-----------------------------------------------------------------------------
// Integral constant helper types <type_traits>
//-----------------------------------------------------------------------------
#if (!defined(_MSC_VER) && (__cplusplus < 201103L)) || (defined(_MSC_VER) && (_MSC_VER < 1500))
/// std::integral_constant
template <typename value_t, value_t V>
struct integral_constant;
/// std::integral_constant
template <typename value_t, value_t V>
struct integral_constant
{
static const value_t value = V;
typedef value_t value_type;
typedef integral_constant<value_t, V> type;
inline __host__ __device__ operator value_type() const
{
return value;
}
inline __host__ __device__ const value_type operator()() const
{
return value;
}
};
#else
using std::integral_constant;
using std::pair;
#endif
/// The type used as a compile-time boolean with true value.
typedef integral_constant<bool, true> true_type;
/// The type used as a compile-time boolean with false value.
typedef integral_constant<bool, false> false_type;
#if (!defined(_MSC_VER) && (__cplusplus < 201402L)) || (defined(_MSC_VER) && (_MSC_VER < 1900))
/// std::bool_constant
template <bool V>
struct bool_constant : nv_std::integral_constant<bool, V>
{};
#else
using std::bool_constant;
#endif
#if (!defined(_MSC_VER) && (__cplusplus < 201103L)) || (defined(_MSC_VER) && (_MSC_VER < 1700))
/// std::nullptr_t
struct nullptr_t {};
#else
using std::nullptr_t;
#endif
//-----------------------------------------------------------------------------
// Conditional metaprogramming <type_traits>
//-----------------------------------------------------------------------------
#if (!defined(_MSC_VER) && (__cplusplus < 201103L)) || (defined(_MSC_VER) && (_MSC_VER < 1600))
/// std::enable_if (true specialization)
template<bool C, typename T = void>
struct enable_if {
typedef T type;
};
/// std::enable_if (false specialization)
template<typename T>
struct enable_if<false, T> { };
/// std::conditional (true specialization)
template<bool B, class T, class F>
struct conditional { typedef T type; };
/// std::conditional (false specialization)
template<class T, class F>
struct conditional<false, T, F> { typedef F type; };
#else
using std::enable_if;
using std::conditional;
#endif
//-----------------------------------------------------------------------------
// Const/volatility specifiers <type_traits>
//-----------------------------------------------------------------------------
#if (!defined(_MSC_VER) && (__cplusplus < 201103L)) || (defined(_MSC_VER) && (_MSC_VER < 1500))
/// std::remove_const (non-const specialization)
template <typename T> struct remove_const { typedef T type; };
/// std::remove_const (const specialization)
template <typename T> struct remove_const<const T> { typedef T type; };
/// std::remove_volatile (non-volatile specialization)
template <typename T> struct remove_volatile { typedef T type; };
/// std::remove_volatile (volatile specialization)
template <typename T> struct remove_volatile<volatile T> { typedef T type; };
/// std::remove_cv
template <typename T>
struct remove_cv {
typedef typename remove_volatile<typename remove_const<T>::type>::type type;
};
#else
using std::remove_const;
using std::remove_volatile;
using std::remove_cv;
#endif
//-----------------------------------------------------------------------------
// Type relationships <type_traits>
//-----------------------------------------------------------------------------
#if (!defined(_MSC_VER) && (__cplusplus < 201103L)) || (defined(_MSC_VER) && (_MSC_VER < 1500))
/// std::is_same (false specialization)
template <typename A, typename B>
struct is_same : false_type
{};
/// std::is_same (true specialization)
template <typename A>
struct is_same<A, A> : true_type
{};
/// Helper for std::is_base_of
template<typename BaseT, typename DerivedT>
struct is_base_of_helper
{
typedef char (&yes)[1];
typedef char (&no)[2];
template<typename B, typename D>
struct dummy
{
operator B*() const;
operator D*();
};
template<typename T>
static yes check(DerivedT*, T);
static no check(BaseT*, int);
static const bool value = sizeof(check(dummy<BaseT, DerivedT>(), int())) == sizeof(yes);
};
/// std::is_base_of
template <typename BaseT, typename DerivedT>
struct is_base_of : integral_constant<
bool,
(is_base_of_helper<typename remove_cv<BaseT>::type, typename remove_cv<DerivedT>::type>::value) ||
(is_same<typename remove_cv<BaseT>::type, typename remove_cv<DerivedT>::type>::value)>
{};
#else
using std::is_same;
using std::is_base_of;
#endif
//-----------------------------------------------------------------------------
// Type properties <type_traits>
//-----------------------------------------------------------------------------
#if (!defined(_MSC_VER) && (__cplusplus < 201103L)) || (defined(_MSC_VER) && (_MSC_VER < 1500))
/// std::is_volatile
template <typename T> struct is_volatile : false_type {};
template <typename T> struct is_volatile<volatile T> : true_type {};
/// Helper for std::is_pointer (false specialization)
template <typename T> struct is_pointer_helper : false_type {};
/// Helper for std::is_pointer (true specialization)
template <typename T> struct is_pointer_helper<T*> : true_type {};
/// std::is_pointer
template <typename T> struct is_pointer : is_pointer_helper<typename remove_cv<T>::type> {};
/// std::is_void
template <typename T>
struct is_void : is_same<void, typename remove_cv<T>::type>
{};
/// std::is_integral
template <typename T> struct is_integral : false_type {};
template <> struct is_integral<char> : true_type {};
template <> struct is_integral<signed char> : true_type {};
template <> struct is_integral<unsigned char> : true_type {};
template <> struct is_integral<short> : true_type {};
template <> struct is_integral<unsigned short> : true_type {};
template <> struct is_integral<int> : true_type {};
template <> struct is_integral<unsigned int> : true_type {};
template <> struct is_integral<long> : true_type {};
template <> struct is_integral<unsigned long> : true_type {};
template <> struct is_integral<long long> : true_type {};
template <> struct is_integral<unsigned long long> : true_type {};
template <typename T> struct is_integral<volatile T> : is_integral<T> {};
template <typename T> struct is_integral<const T> : is_integral<T> {};
template <typename T> struct is_integral<const volatile T> : is_integral<T> {};
/// std::is_floating_point
template <typename T>
struct is_floating_point : integral_constant<
bool,
(is_same<float, typename remove_cv<T>::type>::value ||
is_same<double, typename remove_cv<T>::type>::value)>
{};
/// std::is_arithmetic
template <typename T>
struct is_arithmetic :
integral_constant<bool, (is_integral<T>::value || is_floating_point<T>::value)>
{};
/// std::is_fundamental
template <typename T>
struct is_fundamental : integral_constant<
bool, (is_arithmetic<T>::value ||
is_void<T>::value ||
is_same<nullptr_t, typename remove_cv<T>::type>::value)>
{};
#else
using std::is_volatile;
using std::is_pointer;
using std::is_void;
using std::is_integral;
using std::is_floating_point;
using std::is_arithmetic;
using std::is_fundamental;
#endif
#if (!defined(_MSC_VER) && (__cplusplus < 201103L)) || \
(defined(_MSC_VER) && (_MSC_VER < 1800)) || \
(defined(__GNUG__) && (__GNUC__ < 5))
/**
* std::is_trivially_copyable
*
* This implementation only evaluates true if T is fundamental or pointer
*
* Without help from partial template specializations provided by the user for
* a specific class or struct, this trait will never report that the specified
* class or struct is trivially-copyable ; this is always safe,
* if possibly sub-optimal.
*/
template <typename T>
struct is_trivially_copyable :
integral_constant<bool, (is_fundamental<T>::value || is_pointer<T>::value)>
{};
#else
using std::is_trivially_copyable;
#endif
//-----------------------------------------------------------------------------
// Alignment and layout utilities
//-----------------------------------------------------------------------------
#if (!defined(_MSC_VER) && (__cplusplus < 201103L)) || (defined(_MSC_VER) && (_MSC_VER < 1500))
/// std::alignment_of
template <typename value_t>
struct alignment_of
{
struct pad
{
value_t val;
char byte;
};
enum
{
value = sizeof(pad) - sizeof(value_t)
};
};
#else
template <typename value_t>
struct alignment_of : std::alignment_of<value_t> {};
#endif
/* 16B specializations where 32-bit Win32 host compiler disagrees with device compiler */
template <> struct alignment_of<int4> { enum { value = 16 }; };
template <> struct alignment_of<uint4> { enum { value = 16 }; };
template <> struct alignment_of<float4> { enum { value = 16 }; };
template <> struct alignment_of<long4> { enum { value = 16 }; };
template <> struct alignment_of<ulong4> { enum { value = 16 }; };
template <> struct alignment_of<longlong2> { enum { value = 16 }; };
template <> struct alignment_of<ulonglong2> { enum { value = 16 }; };
template <> struct alignment_of<double2> { enum { value = 16 }; };
template <> struct alignment_of<longlong4> { enum { value = 16 }; };
template <> struct alignment_of<ulonglong4> { enum { value = 16 }; };
template <> struct alignment_of<double4> { enum { value = 16 }; };
// Specializations for volatile/const qualified types
template <typename value_t> struct alignment_of<volatile value_t> : alignment_of<value_t> {};
template <typename value_t> struct alignment_of<const value_t> : alignment_of<value_t> {};
template <typename value_t> struct alignment_of<const volatile value_t> : alignment_of<value_t> {};
#if (!defined(_MSC_VER) && (__cplusplus < 201103L)) || (defined(_MSC_VER) && (_MSC_VER < 1800))
template<size_t Align> struct aligned_chunk;
template<> struct __align__(1) aligned_chunk<1> { uint8_t buff; };
template<> struct __align__(2) aligned_chunk<2> { uint16_t buff; };
template<> struct __align__(4) aligned_chunk<4> { uint32_t buff; };
template<> struct __align__(8) aligned_chunk<8> { uint32_t buff[2]; };
template<> struct __align__(16) aligned_chunk<16> { uint32_t buff[4]; };
template<> struct __align__(32) aligned_chunk<32> { uint32_t buff[8]; };
template<> struct __align__(64) aligned_chunk<64> { uint32_t buff[16]; };
template<> struct __align__(128) aligned_chunk<128> { uint32_t buff[32]; };
template<> struct __align__(256) aligned_chunk<256> { uint32_t buff[64]; };
template<> struct __align__(512) aligned_chunk<512> { uint32_t buff[128]; };
template<> struct __align__(1024) aligned_chunk<1024> { uint32_t buff[256]; };
template<> struct __align__(2048) aligned_chunk<2048> { uint32_t buff[512]; };
template<> struct __align__(4096) aligned_chunk<4096> { uint32_t buff[1024]; };
/// std::aligned_storage
template <size_t Len, size_t Align>
struct aligned_storage
{
typedef aligned_chunk<Align> type[Len / sizeof(aligned_chunk<Align>)];
};
#else
using std::aligned_storage;
#endif
}; // namespace nv_std

View File

@ -1,64 +0,0 @@
/******************************************************************************
* Copyright (c) 2011-2017, NVIDIA CORPORATION. All rights reserved.
*
* Redistribution and use in source and binary forms, with or without
* modification, are not permitted.
*
* THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
* ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
* WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
* DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE FOR ANY
* DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
* (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
* LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
* ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
* (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
* SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
*
******************************************************************************/
#pragma once
/**
* \file
* \brief Pure virtual base class for printable types
*/
#include <iostream>
namespace cutlass {
/******************************************************************************
* printable_t
******************************************************************************/
/**
* Pure virtual base class for printable types
*/
struct printable_t
{
/// Returns the instance as a string
__host__ __device__ inline
virtual char const* to_string() const = 0;
/// Insert the formatted instance into the output stream
virtual void print(std::ostream& out) const = 0;
/// Destructor
virtual ~printable_t() {}
};
/// Insert the formatted \p printable into the output stream
std::ostream& operator<<(
std::ostream& out,
printable_t const& printable)
{
printable.print(out);
return out;
}
} // namespace cutlass

View File

@ -1,74 +0,0 @@
/******************************************************************************
* Copyright (c) 2011-2017, NVIDIA CORPORATION. All rights reserved.
*
* Redistribution and use in source and binary forms, with or without
* modification, are not permitted.
*
* THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
* ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
* WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
* DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE FOR ANY
* DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
* (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
* LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
* ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
* (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
* SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
*
******************************************************************************/
#pragma once
/**
* \file
* \brief Umbrella header file for utilities
*/
#include "debug.h"
#include "device_introspection.h"
#include "io_intrinsics.h"
#include "math.h"
#include "nv_std.h"
#include "printable.h"
#include "matrix_transform.h"
namespace cutlass {
/******************************************************************************
* int_constant
******************************************************************************/
/**
* Shorthand for nv_std::integral_constant of int32_t type
*/
template <int V>
struct int_constant : nv_std::integral_constant<int32_t, V>
{};
/******************************************************************************
* Uninitialized
******************************************************************************/
/**
* \brief A storage-backing wrapper that allows types with non-trivial constructors to be aliased in unions
*/
template <typename T>
struct __align__(16) uninitialized
{
/// Backing storage
uint8_t storage[sizeof(T)];
/// Alias
__host__ __device__ __forceinline__ T& alias()
{
return reinterpret_cast<T&>(*this);
}
};
} // namespace cutlass

View File

@ -1,7 +0,0 @@
/bin/
/gemm-GPU.csv
/gemm-REF.csv
/a.csv
/b.csv
/gp100_schmoo/
/ignore/

View File

@ -1,172 +0,0 @@
#/******************************************************************************
# * Copyright (c) 2011-2017, NVIDIA CORPORATION. All rights reserved.
# *
# * Redistribution and use in source and binary forms, with or without
# * modification, are not permitted.
# *
# * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
# * ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
# * WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
# * DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE FOR ANY
# * DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
# * (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
# * LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
# * ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
# * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
# * SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
# *
#******************************************************************************/
#-------------------------------------------------------------------------------
#
# Makefile usage
#
# make <target> sm=<XX[,YY,ZZ,..]> [transpose=<nn*|nt|tn|tt>] [verbose=<0*|1>] [keep=<0*|1>]
#
# * : default
#
#-------------------------------------------------------------------------------
TEST_DIR := $(dir $(lastword $(MAKEFILE_LIST)))
include ../common.mk
#-------------------------------------------------------------------------------
# Commandline Options
#-------------------------------------------------------------------------------
ifdef transpose
TRANSPOSE := $(transpose)
else
TRANSPOSE := nn
endif
# If defined, GEMMs only compiled with specified alignment restrictions on A and B
# matrices. Otherwise, kernels are compiled for all feasible alignment options, and
# the appropriate kernel is selected.
ifdef alignment
DEFINES += -DGEMM_ALIGNMENT=$(alignment)
endif
# If defined as false, ragged handling can be disabled.
ifdef ragged
DEFINES += -DGEMM_RAGGED=$(ragged)
endif
#-------------------------------------------------------------------------------
# Include and Library paths
#-------------------------------------------------------------------------------
INC += -I$(TEST_DIR)
INC += -I$(BASE_DIR)
LIBS += -lcublas
#-------------------------------------------------------------------------------
# Preprocessor definitions
#-------------------------------------------------------------------------------
ifeq (nt, $(TRANSPOSE))
DEFINES += -DTRANSPOSE_B
else ifeq (tn, $(TRANSPOSE))
DEFINES += -DTRANSPOSE_A
else ifeq (tt, $(TRANSPOSE))
DEFINES += -DTRANSPOSE_A
DEFINES += -DTRANSPOSE_B
endif
NVCCFLAGS += -std=c++11
#-------------------------------------------------------------------------------
# Dependency Lists
#-------------------------------------------------------------------------------
DEPS := $(call rwildcard, $(BASE_DIR),*.h) \
$(call rwildcard, $(BASE_DIR)cgl,*.h) \
$(BASE_DIR)common.mk \
$(TEST_DIR)Makefile
ALL := sgemm \
dgemm \
hgemm \
igemm
#-------------------------------------------------------------------------------
# make default
#-------------------------------------------------------------------------------
default:
#-------------------------------------------------------------------------------
# make clean
#-------------------------------------------------------------------------------
clean :
rm -f bin/*
rm -f *.i* *.cubin *.cu.c *.cudafe* *.fatbin.c *.ptx *.hash *.cu.cpp *.o *.obj* *dlink.* *.res *.fatbin *.module_id
#-------------------------------------------------------------------------------
# make all
#-------------------------------------------------------------------------------
all : $(ALL)
#-------------------------------------------------------------------------------
# make sgemm
#-------------------------------------------------------------------------------
sgemm: bin/sgemm_$(TRANSPOSE)_$(BIN_SUFFIX)
bin/sgemm_$(TRANSPOSE)_$(BIN_SUFFIX) : gemm.cu $(DEPS)
mkdir -p bin
$(NVCC) -DTEST_SGEMM $(DEFINES) $(SM_TARGETS) -o $@ gemm.cu $(NVCCFLAGS) $(CPU_ARCH) $(INC) $(LIBINC) $(LIBS)
#-------------------------------------------------------------------------------
# make dgemm
#-------------------------------------------------------------------------------
dgemm: bin/dgemm_$(TRANSPOSE)_$(BIN_SUFFIX)
bin/dgemm_$(TRANSPOSE)_$(BIN_SUFFIX) : gemm.cu $(DEPS)
mkdir -p bin
$(NVCC) -DTEST_DGEMM $(DEFINES) $(SM_TARGETS) -o $@ gemm.cu $(NVCCFLAGS) $(CPU_ARCH) $(INC) $(LIBINC) $(LIBS)
#-------------------------------------------------------------------------------
# make hgemm
#-------------------------------------------------------------------------------
hgemm: bin/hgemm_$(TRANSPOSE)_$(BIN_SUFFIX)
bin/hgemm_$(TRANSPOSE)_$(BIN_SUFFIX) : gemm.cu $(DEPS)
mkdir -p bin
$(NVCC) -DTEST_HGEMM $(DEFINES) $(SM_TARGETS) -o $@ gemm.cu $(NVCCFLAGS) $(CPU_ARCH) $(INC) $(LIBINC) $(LIBS)
#-------------------------------------------------------------------------------
# make igemm
#-------------------------------------------------------------------------------
igemm: bin/igemm_$(TRANSPOSE)_$(BIN_SUFFIX)
bin/igemm_$(TRANSPOSE)_$(BIN_SUFFIX) : gemm.cu $(DEPS)
mkdir -p bin
$(NVCC) -DTEST_IGEMM $(DEFINES) $(SM_TARGETS) -o $@ gemm.cu $(NVCCFLAGS) $(CPU_ARCH) $(INC) $(LIBINC) $(LIBS)
#-------------------------------------------------------------------------------
# make wgemm
#-------------------------------------------------------------------------------
wgemm: bin/wgemm_$(TRANSPOSE)_$(BIN_SUFFIX)
bin/wgemm_$(TRANSPOSE)_$(BIN_SUFFIX) : gemm.cu $(DEPS)
mkdir -p bin
$(NVCC) -DTEST_WGEMM -DWMMA $(DEFINES) $(SM_TARGETS) -o $@ gemm.cu $(NVCCFLAGS) $(CPU_ARCH) $(INC) $(LIBINC) $(LIBS)

View File

@ -1,292 +0,0 @@
/******************************************************************************
* Copyright (c) 2011-2017, NVIDIA CORPORATION. All rights reserved.
*
* Redistribution and use in source and binary forms, with or without
* modification, are not permitted.
*
* THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
* ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
* WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
* DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE FOR ANY
* DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
* (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
* LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
* ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
* (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
* SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
*
******************************************************************************/
#pragma once
/**
* \file
* C++ interface for dispatching CUBLAS GEMM calls
*/
#include <cublas_v2.h>
namespace cutlass {
/******************************************************************************
* cuBLAS dispatch entrypoints
******************************************************************************/
/**
* Dispatch cuBLAS igemm
*/
cublasStatus_t cublas_gemm_dispatch(
cublasHandle_t cublas_handle, ///< CUBLAS handle
cublasOperation_t transform_a, ///< Transform op(A) that is non- or (conj.) transpose.
cublasOperation_t transform_b, ///< Transform op(B) that is non- or (conj.) transpose.
int m, ///< Height in rows of op(A) and C
int n, ///< Width in columns of op(B) and C
int k, ///< Width in columns of op(A) and height in rows of op(B)
int32_t alpha, ///< Scalar used for multiplicands
int8_t *d_a, ///< Device pointer to matrix A array values
int8_t *d_b, ///< Device pointer to matrix B array values
int32_t beta, ///< Scalar used for addend
int32_t *d_c, ///< Device pointer to matrix C array values
cudaStream_t stream = 0, ///< CUDA stream to launch kernels within. Default is stream<sub>0</sub>.
bool debug_synchronous = false) ///< Whether or not to synchronize the stream after every kernel launch to check for errors.
{
return cublasGemmEx(
cublas_handle,
transform_a,
transform_b,
m,
n,
k,
(void*) &alpha,
(void*) d_a,
CUDA_R_8I,
(transform_a == CUBLAS_OP_N) ? m : k,
(void*) d_b,
CUDA_R_8I,
(transform_b == CUBLAS_OP_N) ? k : n,
(void*) &beta,
(void*) d_c,
CUDA_R_32I,
m,
CUDA_R_32I,
CUBLAS_GEMM_DFALT);
}
/**
* Dispatch cuBLAS hgemm
*/
cublasStatus_t cublas_gemm_dispatch(
cublasHandle_t cublas_handle, ///< CUBLAS handle
cublasOperation_t transform_a, ///< Transform op(A) that is non- or (conj.) transpose.
cublasOperation_t transform_b, ///< Transform op(B) that is non- or (conj.) transpose.
int m, ///< Height in rows of op(A) and C
int n, ///< Width in columns of op(B) and C
int k, ///< Width in columns of op(A) and height in rows of op(B)
__half alpha, ///< Scalar used for multiplicands
__half *d_a, ///< Device pointer to matrix A array values
__half *d_b, ///< Device pointer to matrix B array values
__half beta, ///< Scalar used for addend
__half *d_c, ///< Device pointer to matrix C array values
cudaStream_t stream = 0, ///< CUDA stream to launch kernels within. Default is stream<sub>0</sub>.
bool debug_synchronous = false) ///< Whether or not to synchronize the stream after every kernel launch to check for errors.
{
return cublasHgemm(
cublas_handle, transform_a, transform_b,
m, n, k,
&alpha,
d_a,
(transform_a == CUBLAS_OP_N) ? m : k,
d_b,
(transform_b == CUBLAS_OP_N) ? k : n,
&beta,
d_c,
m);
}
/**
* Dispatch cuBLAS sgemm
*/
cublasStatus_t cublas_gemm_dispatch(
cublasHandle_t cublas_handle, ///< CUBLAS handle
cublasOperation_t transform_a, ///< Transform op(A) that is non- or (conj.) transpose.
cublasOperation_t transform_b, ///< Transform op(B) that is non- or (conj.) transpose.
int m, ///< Height in rows of op(A) and C
int n, ///< Width in columns of op(B) and C
int k, ///< Width in columns of op(A) and height in rows of op(B)
float alpha, ///< Scalar used for multiplicands
float *d_a, ///< Device pointer to matrix A array values
float *d_b, ///< Device pointer to matrix B array values
float beta, ///< Scalar used for addend
float *d_c, ///< Device pointer to matrix C array values
cudaStream_t stream = 0, ///< CUDA stream to launch kernels within. Default is stream<sub>0</sub>.
bool debug_synchronous = false) ///< Whether or not to synchronize the stream after every kernel launch to check for errors.
{
return cublasSgemm(
cublas_handle, transform_a, transform_b,
m, n, k,
&alpha,
d_a,
(transform_a == CUBLAS_OP_N) ? m : k,
d_b,
(transform_b == CUBLAS_OP_N) ? k : n,
&beta,
d_c,
m);
}
/**
* Dispatch cuBLAS dgemm
*/
cublasStatus_t cublas_gemm_dispatch(
cublasHandle_t cublas_handle, ///< CUBLAS handle
cublasOperation_t transform_a, ///< Transform op(A) that is non- or (conj.) transpose.
cublasOperation_t transform_b, ///< Transform op(B) that is non- or (conj.) transpose.
int m, ///< Height in rows of op(A) and C
int n, ///< Width in columns of op(B) and C
int k, ///< Width in columns of op(A) and height in rows of op(B)
double alpha, ///< Scalar used for multiplicands
double *d_a, ///< Device pointer to matrix A array values
double *d_b, ///< Device pointer to matrix B array values
double beta, ///< Scalar used for addend
double *d_c, ///< Device pointer to matrix C array values
cudaStream_t stream = 0, ///< CUDA stream to launch kernels within. Default is stream<sub>0</sub>.
bool debug_synchronous = false) ///< Whether or not to synchronize the stream after every kernel launch to check for errors.
{
return cublasDgemm(
cublas_handle, transform_a, transform_b,
m, n, k,
&alpha,
d_a, (transform_a == CUBLAS_OP_N) ? m : k,
d_b, (transform_b == CUBLAS_OP_N) ? k : n,
&beta,
d_c, m);
}
/**
* Dispatch cuBLAS Tensor Cores GEMM
*/
cublasStatus_t cublas_gemm_dispatch(
cublasHandle_t cublas_handle, ///< CUBLAS handle
cublasOperation_t transform_a, ///< Transform op(A) that is non- or (conj.) transpose.
cublasOperation_t transform_b, ///< Transform op(B) that is non- or (conj.) transpose.
int m, ///< Height in rows of op(A) and C
int n, ///< Width in columns of op(B) and C
int k, ///< Width in columns of op(A) and height in rows of op(B)
float alpha, ///< Scalar used for multiplicands
half *d_a, ///< Device pointer to matrix A array values
half *d_b, ///< Device pointer to matrix B array values
float beta, ///< Scalar used for addend
float *d_c, ///< Device pointer to matrix C array values
cudaStream_t stream = 0, ///< CUDA stream to launch kernels within. Default is stream<sub>0</sub>.
bool debug_synchronous = false) ///< Whether or not to synchronize the stream after every kernel launch to check for errors.
{
return cublasGemmEx(
cublas_handle,
transform_a,
transform_b,
m,
n,
k,
(void*) &alpha,
(void*) d_a,
CUDA_R_16F,
(transform_a == CUBLAS_OP_N) ? m : k,
(void*) d_b,
CUDA_R_16F,
(transform_b == CUBLAS_OP_N) ? k : n,
(void*) &beta,
(void*) d_c,
CUDA_R_32F,
m,
CUDA_R_32F,
CUBLAS_GEMM_DFALT_TENSOR_OP);
}
/**
* Uses cuBLAS to compute gemm on device matrices (unspecialized)
*/
template <
gemm::tiling_strategy::kind_t _TilingStrategy, ///< Tile-sizing classification category
math_operation_class_t _math_op,
matrix_transform_t::kind_t _TransformA, ///< Transformation op for matrix A
matrix_transform_t::kind_t _TransformB, ///< Transformation op for matrix B
typename _value, ///< Multiplicand value type (matrices A and B)
typename _accum ///< Accumulator value type (matrix C and scalars)
>
struct cublas_gemm
{
//
// Type alias definitions
//
static const gemm::tiling_strategy::kind_t TilingStrategy = _TilingStrategy;
static const math_operation_class_t math_op = _math_op;
static const matrix_transform_t::kind_t TransformA = _TransformA;
static const matrix_transform_t::kind_t TransformB = _TransformB;
using value_t = _value;
using accum_t = _accum;
/// Launches a GEMM
gemm::launch_configuration operator()(
cublasHandle_t cublas_handle, ///< CUBLAS handle
int m,
int n,
int k,
value_t *A, ///< A matrix
value_t *B, ///< B matrix
accum_t *C, ///< C matrix
accum_t alpha, ///< Scalar used for multiplicands
accum_t beta, ///< Scalar used for addend
cudaStream_t stream = 0, ///< CUDA stream to launch kernels within. Default is stream<sub>0</sub>.
bool debug_synchronous = false) ///< Whether or not to synchronize the stream after every kernel launch to check for errors.
{
cublasStatus_t cublas_error = cublas_gemm_dispatch(
cublas_handle,
(cublasOperation_t) TransformA,
(cublasOperation_t) TransformB,
m,
n,
k,
alpha,
A,
B,
beta,
C,
stream,
debug_synchronous);
cudaError_t error;
if (cublas_error != CUBLAS_STATUS_SUCCESS)
{
if (cublas_error == CUBLAS_STATUS_NOT_SUPPORTED) {
return gemm::launch_configuration(cudaErrorInvalidValue);
}
error = cudaGetLastError();
if (error == cudaSuccess) {
return gemm::launch_configuration(cudaErrorUnknown);
}
return error;
}
// Check for failure to launch
if (CUDA_PERROR_DEBUG(error = cudaPeekAtLastError()))
return gemm::launch_configuration(error);
// Sync the stream if specified to flush runtime errors
if (debug_synchronous && (CUDA_PERROR_DEBUG(error = cudaStreamSynchronize(stream))))
return gemm::launch_configuration(error);
return gemm::launch_configuration(error);
}
};
} // namespace cutlass

View File

@ -1,253 +0,0 @@
/******************************************************************************
* Copyright (c) 2011-2017, NVIDIA CORPORATION. All rights reserved.
*
* Redistribution and use in source and binary forms, with or without
* modification, are not permitted.
*
* THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
* ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
* WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
* DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE FOR ANY
* DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
* (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
* LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
* ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
* (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
* SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
*
******************************************************************************/
#pragma once
/**
* \file Dispatch routines for CUTLASS GEMM kernels
*/
// CUDA includes
#include <cublas_v2.h>
// Cutlass GEMM API
#include <cutlass/util/util.h>
#include <cutlass/gemm/dispatch.h>
#include <cutlass/gemm/epilogue_function.h>
// Test utilities
#include "util/type_conversion.h"
namespace cutlass {
/******************************************************************************
* Cutlass dispatch entrypoints
******************************************************************************/
//
// Compile-time overrides for alignment and ragged handling.
//
// If zero, all feasible alignment options are supported.
#ifndef GEMM_ALIGNMENT
#define GEMM_ALIGNMENT 0
#endif
// If true, kernels are compiled with ragged handling enabled.
#ifndef GEMM_RAGGED
#define GEMM_RAGGED true
#endif
//
// Dispatch logic given problem size specialization, math operation class, layout
// and type of operands, and epilogue operation.
//
/**
* Cutlass GEMM dispatch
*/
template <
gemm::tiling_strategy::kind_t _TilingStrategy, ///< Tile-sizing classification category
math_operation_class_t _math_op, // Indicates
matrix_transform_t::kind_t _TransformA, ///< Transformation op for matrix A
matrix_transform_t::kind_t _TransformB, ///< Transformation op for matrix B
typename _value, ///< Multiplicand value type (matrices A and B)
typename _accum, ///< Accumulator value type (matrix C and scalars)
typename _epilogue_op_t ///< Epilogue opeartion to update matrix C
= gemm::blas_scaled_epilogue<_accum, _accum, _accum>
>
struct cutlass_gemm_dispatch
{
//
// Type alias definitions
//
static const gemm::tiling_strategy::kind_t TilingStrategy = _TilingStrategy;
static const math_operation_class_t math_op = _math_op;
static const matrix_transform_t::kind_t TransformA = _TransformA;
static const matrix_transform_t::kind_t TransformB = _TransformB;
using value_t = _value;
using accum_t = _accum;
using epilogue_op_t = _epilogue_op_t;
//
// Methods
//
/// Returns leading dimension for A matrix operand
int leading_dim_a(int m, int k) const
{
return (TransformA == matrix_transform_t::NonTranspose ? m : k);
}
/// Returns leading dimension for B matrix operand
int leading_dim_b(int k, int n) const
{
return (TransformB == matrix_transform_t::NonTranspose ? k : n);
}
/// Launches a GEMM
template <int operand_alignment, int accumulator_alignment>
gemm::launch_configuration launch(
int m,
int n,
int k,
epilogue_op_t epilogue_op,
value_t *A,
value_t *B,
accum_t *C,
cudaStream_t stream = 0,
bool debug_synchronous = false)
{
return gemm::device_gemm<
TilingStrategy,
math_op,
TransformA,
operand_alignment,
TransformB,
operand_alignment,
value_t,
accum_t,
epilogue_op_t,
accumulator_alignment>
(
m,
n,
k,
epilogue_op,
A,
B,
C,
stream,
debug_synchronous);
}
/// Dispatches a CUTLASS GEMM
gemm::launch_configuration operator()(
cublasHandle_t handle, ///< CUBLAS handle
int m, ///< Rows of GEMM problem
int n, ///< Columns of GEMM problem
int k, ///< Inner dimension of GEMM problem
value_t *A, ///< A matrix
value_t *B, ///< B matrix
accum_t *C, ///< C matrix
accum_t alpha, ///< Scalar used for multiplicands
accum_t beta, ///< Scalar used for addend
cudaStream_t stream = 0, ///< CUDA stream to launch kernels within.
bool debug_synchronous = false) ///< Whether or not to synchronize the stream
/// after every kernel launch to check for errors.
{
// Forces kernel selection to choose specific alignment (in bytes)
int const force_operand_alignment = GEMM_ALIGNMENT;
// Problem size must be multiple of the smallest vector load size
typedef value_t operand_load_t;
int const accumulator_alignment = sizeof(accum_t);
int const lda = leading_dim_a(m, k);
int const ldb = leading_dim_b(k, n);
epilogue_op_t epilogue(alpha, beta);
// TODO: opportunity for metaprogramming loop
// Prefer the largest granularity of vector load that is compatible with
// problem size and data alignment.
if ((!force_operand_alignment || force_operand_alignment == 16) &&
!((sizeof(operand_load_t) * lda) % 16) &&
!((sizeof(operand_load_t) * ldb) % 16))
{
#if !(GEMM_ALIGNMENT) || (GEMM_ALIGNMENT == 16)
return launch<__NV_STD_MAX(16, sizeof(value_t)), accumulator_alignment>(
m,
n,
k,
epilogue,
A,
B,
C,
stream,
debug_synchronous);
#endif
}
else if ((!force_operand_alignment || force_operand_alignment == 8) &&
!((sizeof(operand_load_t) * lda) % 8) &&
!((sizeof(operand_load_t) * ldb) % 8))
{
#if !(GEMM_ALIGNMENT) || (GEMM_ALIGNMENT == 8)
return launch<__NV_STD_MAX(8, sizeof(value_t)), accumulator_alignment>(
m,
n,
k,
epilogue,
A,
B,
C,
stream,
debug_synchronous);
#endif
}
else if ((!force_operand_alignment || force_operand_alignment == 4) &&
!((sizeof(operand_load_t) * lda) % 4) &&
!((sizeof(operand_load_t) * ldb) % 4))
{
#if !(GEMM_ALIGNMENT) || (GEMM_ALIGNMENT == 4)
return launch<__NV_STD_MAX(4, sizeof(value_t)), accumulator_alignment>(
m,
n,
k,
epilogue,
A,
B,
C,
stream,
debug_synchronous);
#endif
}
else if ((!force_operand_alignment || force_operand_alignment == 2) &&
!((sizeof(operand_load_t) * lda) % 2) &&
!((sizeof(operand_load_t) * ldb) % 2))
{
// 16-bit alignment only supported for HGEMM
#if defined(TEST_HGEMM) || defined(TEST_WGEMM)
#if !(GEMM_ALIGNMENT) || (GEMM_ALIGNMENT == 2)
return launch<__NV_STD_MAX(2, sizeof(value_t)), accumulator_alignment>(
m,
n,
k,
epilogue,
A,
B,
C,
stream,
debug_synchronous);
#endif
#endif
}
return gemm::launch_configuration(cudaErrorInvalidValue);
}
};
} // namespace cutlass

View File

@ -1,564 +0,0 @@
/******************************************************************************
* Copyright (c) 2011-2017, NVIDIA CORPORATION. All rights reserved.
*
* Redistribution and use in source and binary forms, with or without
* modification, are not permitted.
*
* THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
* ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
* WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
* DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE FOR ANY
* DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
* (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
* LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
* ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
* (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
* SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
*
******************************************************************************/
/**
* \file gemm.cu
* GEMM test driver
*
*/
#include <iostream>
#include <typeinfo>
#include <random>
#include <stdint.h>
// CUBLAS GEMM API
#include <cublas_v2.h>
// Set Cutlass debug macro to enable console printing of library errors
#define DEBUG
#if defined(WMMA)
// Conditionally include WMMA headers (CUDA 9 Preview Feature)
#include <mma.h>
#endif
// Cutlass GEMM API
#include <cutlass/util/util.h>
#include <cutlass/gemm/dispatch.h>
#include <cutlass/gemm/epilogue_function.h>
// Test utilities
#include "util/command_line.h"
#include "util/half.h"
#include "util/matrix.h"
#include "util/timer.h"
#include "util/type_conversion.h"
// Dispatch routines to CUBLAS and CUTLASS
#include "cublas_dispatch.h"
#include "cutlass_dispatch.h"
/******************************************************************************
* Globals, constants and typedefs
******************************************************************************/
using namespace cutlass;
/// CUBLAS handle
cublasHandle_t g_cublas_handle;
/// The device-id of the current device
int g_device_id = -1;
/// The number of timing iterations to invoke
int g_timing_iterations = -1;
/// The number of randomly-sized problems to schmoo
int g_schmoo = 0;
/******************************************************************************
* Number generation
******************************************************************************/
/**
* Simple low-integer generator
*/
struct simple_gen
{
std::default_random_engine generator;
std::uniform_int_distribution<int> distribution;
/// Constructor
simple_gen(int max) : distribution(max * -1, max)
{}
/// Functor
int operator()()
{
return distribution(generator);
}
};
/******************************************************************************
* Test execution
******************************************************************************/
/**
* Compute C = (alpha * A * B) + (beta * C)
*/
template <
typename test_func_t, ///< Test function type
matrix_transform_t::kind_t TransformA, ///< Transformation op for matrix A
matrix_transform_t::kind_t TransformB, ///< Transformation op for matrix B
typename value_t, ///< Multiplicand value type (matrices A and B)
typename accum_t> ///< Accumulator value type (matrix C and scalars)
bool test(
int m, ///< Height of C in rows
int n, ///< Width of C in columns
int k, ///< Width (height) of A (B)
accum_t alpha, ///< Multiplicand scalar
accum_t beta) ///< Addend scalar
{
cudaStream_t stream = 0;
//
// Initialize matrices
//
matrix<value_t> A(
(TransformA == matrix_transform_t::NonTranspose) ? m : k,
(TransformA == matrix_transform_t::NonTranspose) ? k : m);
matrix<value_t> B(
(TransformB == matrix_transform_t::NonTranspose) ? k : n,
(TransformB == matrix_transform_t::NonTranspose) ? n : k);
matrix<accum_t> C(m, n);
// initialized matrices with small values precisely representable as integers
simple_gen a_gen(3);
simple_gen b_gen(5);
A.fill_random(a_gen);
B.fill_random(b_gen);
C.fill_ramp(0,0);
// // Alternatively, initialize with procedural values to simplify debugging incorrect results
// A.fill_ramp(1,2);
// B.fill_ramp(1,1);
// Sync to device
A.sync_device();
B.sync_device();
C.sync_device();
CUDA_PERROR(cudaPeekAtLastError());
CUDA_PERROR(cudaDeviceSynchronize());
//
// Run test once with debug-synchronous enabled and check result
//
if (!g_schmoo) printf("\n");
test_func_t test_func;
C.fill_ramp(0, 0);
C.sync_device();
cudaError_t error = test_func(
g_cublas_handle,
m,
n,
k,
A.d_data(),
B.d_data(),
C.d_data(),
alpha,
beta,
stream,
!g_schmoo).result;
bool not_applicable = (error == cudaErrorInvalidValue);
bool is_failed = false;
if (not_applicable)
{
printf(", NA");
}
else
{
CUDA_PERROR(error);
// Compute reference check if wont take too long on CPU
if ((!g_schmoo) && (m * n <= 1024 * 1024))
{
matrix<accum_t> ref_C(m, n);
ref_C.fill_ramp(0, 0);
ref_C.gemm(TransformA, TransformB, alpha, A, B, beta);
C.sync_host();
is_failed = (C != ref_C);
if (!g_schmoo)
{
if (is_failed)
{
printf("FAIL, ");
std::ofstream file_a("a.csv");
A.write_matrix(file_a);
std::ofstream file_b("b.csv");
B.write_matrix(file_b);
std::ofstream file_d("gemm-REF.csv");
ref_C.write_matrix(file_d);
std::ofstream file_c("gemm-GPU.csv");
C.write_matrix(file_c);
}
else
{
printf("PASS, ");
}
}
}
fflush(stdout);
//
// Warmup and timing iterations
//
if (g_timing_iterations > 0)
{
// Warmup for 1/100 of the timing iterations (minimum of 2)
for (int i = 0; i < __NV_STD_MAX(2, (g_timing_iterations + 99) / 100); ++i)
{
CUDA_PERROR(test_func(
g_cublas_handle,
m,
n,
k,
A.d_data(),
B.d_data(),
C.d_data(),
alpha,
beta,
stream,
false).result);
}
}
// Conduct timing iterations
double elapsed_ms = 0;
gpu_timer timer;
timer.start();
for (int i = 0; i < g_timing_iterations; i++)
{
CUDA_PERROR(test_func(
g_cublas_handle,
m,
n,
k,
A.d_data(),
B.d_data(),
C.d_data(),
alpha,
beta,
stream,
false).result);
}
timer.stop();
elapsed_ms += timer.elapsed_millis();
double avg_ms = elapsed_ms / g_timing_iterations;
// Display performance
if (g_timing_iterations > 0)
{
int64_t num_flops = (2 * int64_t(m) * int64_t(n) * int64_t(k)) + (2 * int64_t(m) * int64_t(n));
double gflops_per_sec = double(num_flops) / avg_ms / 1.0e6;
if (g_schmoo)
{
if (is_failed)
printf("F");
printf(", %.3f", gflops_per_sec);
// Sleep for a few milliseconds to cool
sleep_millis(10);
}
else
{
printf("Avg runtime: %.3f ms, total flops: %lld, GFLOP/s: %.2f\n",
avg_ms,
num_flops,
gflops_per_sec);
}
fflush(stdout);
}
}
return is_failed;
}
/**
* Compute C = (alpha * A * B) + (beta * C)
*/
template <
math_operation_class_t math_op,
matrix_transform_t::kind_t TransformA, ///< Transformation op for matrix A
matrix_transform_t::kind_t TransformB, ///< Transformation op for matrix B
typename value_t, ///< Multiplicand value type (matrices A and B)
typename accum_t> ///< Accumulator value type (matrix C and scalars)
bool test(
int m, ///< Height of C in rows
int n, ///< Width of C in columns
int k, ///< Width (height) of A (B)
accum_t alpha, ///< Multiplicand scalar
accum_t beta) ///< Addend scalar
{
uint64_t flop_base = 1ull << 41;
int max_timing_iterations = 10000;
int min_timing_iterations = 10;
bool test_error = false;
// Scale the number of timing iterations with respect to problem size (if not specified on commandline)
if ((g_timing_iterations < 0) || g_schmoo)
{
uint64_t num_flops = (2 * uint64_t(m) * uint64_t(n) * uint64_t(k)) + (2 * uint64_t(m) * uint64_t(n));
g_timing_iterations = (int) ((flop_base / sizeof(value_t)) / num_flops);
g_timing_iterations = (int) __NV_STD_MIN(max_timing_iterations, g_timing_iterations);
g_timing_iterations = (int) __NV_STD_MAX(min_timing_iterations, g_timing_iterations);
}
if (g_schmoo)
{
printf("%d, %d, %d, %c%c, %d, %d",
m, n, k,
(TransformA == matrix_transform_t::NonTranspose) ? 'n' : 't',
(TransformB == matrix_transform_t::NonTranspose) ? 'n' : 't',
m * n,
g_timing_iterations);
}
else
{
printf("\n------------------------------------------------------------\n");
printf("%dx%dx%d, GEMM_%c%c, %d C elements, %d timing iterations\n",
m, n, k,
(TransformA == matrix_transform_t::NonTranspose) ? 'n' : 't',
(TransformB == matrix_transform_t::NonTranspose) ? 'n' : 't',
m * n,
g_timing_iterations);
}
fflush(stdout);
// CUBLAS
test_error |= test<
cublas_gemm<gemm::tiling_strategy::Unknown, math_op, TransformA, TransformB, value_t, accum_t>,
TransformA,
TransformB,
value_t,
accum_t>(m, n, k, accum_t(alpha), accum_t(beta));
// CUTLASS
test_error |= test<
cutlass_gemm_dispatch<gemm::tiling_strategy::Small, math_op, TransformA, TransformB, value_t, accum_t>,
TransformA,
TransformB,
value_t,
accum_t>(m, n, k, accum_t(alpha), accum_t(beta));
test_error |= test<
cutlass_gemm_dispatch<gemm::tiling_strategy::Medium, math_op, TransformA, TransformB, value_t, accum_t>,
TransformA,
TransformB,
value_t,
accum_t>(m, n, k, accum_t(alpha), accum_t(beta));
test_error |= test<
cutlass_gemm_dispatch<gemm::tiling_strategy::Large, math_op, TransformA, TransformB, value_t, accum_t>,
TransformA,
TransformB,
value_t,
accum_t>(m, n, k, accum_t(alpha), accum_t(beta));
test_error |= test<
cutlass_gemm_dispatch<gemm::tiling_strategy::Tall, math_op, TransformA, TransformB, value_t, accum_t>,
TransformA,
TransformB,
value_t,
accum_t>(m, n, k, accum_t(alpha), accum_t(beta));
test_error |= test<
cutlass_gemm_dispatch<gemm::tiling_strategy::Wide, math_op, TransformA, TransformB, value_t, accum_t>,
TransformA,
TransformB,
value_t,
accum_t>(m, n, k, accum_t(alpha), accum_t(beta));
test_error |= test<
cutlass_gemm_dispatch<gemm::tiling_strategy::Huge, math_op, TransformA, TransformB, value_t, accum_t>,
TransformA,
TransformB,
value_t,
accum_t>(m, n, k, accum_t(alpha), accum_t(beta));
return test_error;
}
/******************************************************************************
* Main
******************************************************************************/
/**
* Main
*/
int main(int argc, const char **argv)
{
//
// Problem type (compiler-supplied so we don't compile everything)
//
// Define value_t and accum_t (multiplicand and accumulator types, respectively)
#if defined(TEST_SGEMM)
typedef float value_t;
typedef float accum_t;
const math_operation_class_t math_op = math_operation_class_t::scalar;
#elif defined(TEST_DGEMM)
typedef double value_t;
typedef double accum_t;
const math_operation_class_t math_op = math_operation_class_t::scalar;
#elif defined(TEST_HGEMM)
typedef __half value_t;
typedef __half accum_t;
const math_operation_class_t math_op = math_operation_class_t::scalar;
#elif defined(TEST_IGEMM)
typedef int8_t value_t;
typedef int32_t accum_t;
const math_operation_class_t math_op = math_operation_class_t::scalar;
#elif defined(TEST_WGEMM)
typedef half value_t;
typedef float accum_t;
const math_operation_class_t math_op = math_operation_class_t::matrix;
#else
#error Unknown GEMM type requested.
#endif
// Define transpose constants
#ifdef TRANSPOSE_A
static const matrix_transform_t::kind_t TransformA = matrix_transform_t::Transpose;
#else
static const matrix_transform_t::kind_t TransformA = matrix_transform_t::NonTranspose;
#endif
#ifdef TRANSPOSE_B
static const matrix_transform_t::kind_t TransformB = matrix_transform_t::Transpose;
#else
static const matrix_transform_t::kind_t TransformB = matrix_transform_t::NonTranspose;
#endif
//
// Commandline parsing
//
// Initialize command line
command_line args(argc, argv);
int m_factor = args.device_prop.multiProcessorCount * 128;
int m = round_nearest(4096, m_factor);
int k = 4096;
int n = 4096;
float alpha = 1.0;
float beta = 0.0;
g_device_id = args.device_id;
args.get_cmd_line_argument("m", m);
args.get_cmd_line_argument("n", n);
args.get_cmd_line_argument("k", k);
args.get_cmd_line_argument("i", g_timing_iterations);
args.get_cmd_line_argument("alpha", alpha);
args.get_cmd_line_argument("beta", beta);
args.get_cmd_line_argument("schmoo", g_schmoo);
// Print usage
if (args.check_cmd_line_flag("help"))
{
printf("%s "
"[--help] "
"[--i=<timing iterations>] "
"[--device=<device-id>] "
"[--alpha=<alpha> --beta=<beta>] "
"[--schmoo=<samples> || --m=<height> --n=<width> --k=<depth>]"
"\n", argv[0]);
exit(0);
}
// Initialize cuBLAS
if (cublasCreate(&g_cublas_handle) != CUBLAS_STATUS_SUCCESS)
{
fprintf(stderr, "cublasCreate() failed\n");
exit(1);
}
bool test_error = false;
if (g_schmoo)
{
// Run a schmoo of problem sizes
printf("M, N, K, transpose, total_flops, timing_iterations, sol_flop/s, cublas_sol, cutlass_small_sol, cutlass_med_sol, cutlass_large_sol, cutlass_tall_sol, cutlass_wide_sol, cutlass_huge_sol\n");
// Generate power-law distribution from [32, 16384)
std::mt19937 gen(0);
std::uniform_real_distribution<float> dis(5, 14);
for (int i = 0; i < g_schmoo; ++i)
{
int m = int(pow(float(2), dis(gen)));
int n = int(pow(float(2), dis(gen)));
int k = int(pow(float(2), dis(gen)));
// Round m and n to nearest multiple of 32 if < 128, otherwise to the nearest 128
m = (m < 128) ?
round_nearest(m, 32) :
round_nearest(m, 128);
n = (n < 128) ?
round_nearest(n, 32) :
round_nearest(n, 128);
// Round k to the nearest 16
k = (sizeof(value_t) == 1) ?
round_nearest(k, 32) :
round_nearest(k, 16);
test_error |= test<math_op, TransformA, TransformB, value_t, accum_t>(
m, n, k,
from_float<accum_t>(alpha),
from_float<accum_t>(beta));
printf("\n"); fflush(stdout);
}
}
else
{
// Test a single GEMM problem size
test_error |= test<math_op, TransformA, TransformB, value_t, accum_t>(
m,
n,
k,
from_float<accum_t>(alpha),
from_float<accum_t>(beta));
}
// Cleanup
cublasDestroy(g_cublas_handle);
return test_error;
}

View File

@ -1,312 +0,0 @@
/******************************************************************************
* Copyright (c) 2011-2017, NVIDIA CORPORATION. All rights reserved.
*
* Redistribution and use in source and binary forms, with or without
* modification, are not permitted.
*
* THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
* ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
* WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
* DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE FOR ANY
* DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
* (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
* LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
* ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
* (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
* SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
*
******************************************************************************/
#pragma once
/**
* \file
* Utility for parsing command line arguments
*/
#include <string>
#include <vector>
#include <sstream>
#include <iostream>
#include <limits>
#include <cuda_runtime.h>
#include <cutlass/util/debug.h>
namespace cutlass {
/******************************************************************************
* command_line
******************************************************************************/
/**
* Utility for parsing command line arguments
*/
struct command_line
{
std::vector<std::string> keys;
std::vector<std::string> values;
std::vector<std::string> args;
int device_id;
cudaDeviceProp device_prop;
float device_giga_bandwidth;
size_t device_free_physmem;
size_t device_total_physmem;
/**
* Constructor
*/
command_line(int argc, const char **argv, int device_id = -1) :
keys(10),
values(10),
device_id(device_id)
{
using namespace std;
for (int i = 1; i < argc; i++)
{
string arg = argv[i];
if ((arg[0] != '-') || (arg[1] != '-'))
{
args.push_back(arg);
continue;
}
string::size_type pos;
string key, val;
if ((pos = arg.find('=')) == string::npos) {
key = string(arg, 2, arg.length() - 2);
val = "";
} else {
key = string(arg, 2, pos - 2);
val = string(arg, pos + 1, arg.length() - 1);
}
keys.push_back(key);
values.push_back(val);
}
// Initialize device
CUDA_PERROR_EXIT(device_init());
}
/**
* Checks whether a flag "--<flag>" is present in the commandline
*/
bool check_cmd_line_flag(const char* arg_name)
{
using namespace std;
for (int i = 0; i < int(keys.size()); ++i)
{
if (keys[i] == string(arg_name))
return true;
}
return false;
}
/**
* Returns number of naked (non-flag and non-key-value) commandline parameters
*/
template <typename value_t>
int num_naked_args()
{
return args.size();
}
/**
* Returns the commandline parameter for a given index (not including flags)
*/
template <typename value_t>
void get_cmd_line_argument(int index, value_t &val)
{
using namespace std;
if (index < args.size()) {
istringstream str_stream(args[index]);
str_stream >> val;
}
}
/**
* Returns the value specified for a given commandline parameter --<flag>=<value>
*/
template <typename value_t>
void get_cmd_line_argument(const char *arg_name, value_t &val)
{
using namespace std;
for (int i = 0; i < int(keys.size()); ++i)
{
if (keys[i] == string(arg_name))
{
istringstream str_stream(values[i]);
str_stream >> val;
}
}
}
/**
* Returns the values specified for a given commandline parameter --<flag>=<value>,<value>*
*/
template <typename value_t>
void get_cmd_line_arguments(
const char *arg_name,
std::vector<value_t> &vals,
char sep = ',')
{
using namespace std;
if (check_cmd_line_flag(arg_name))
{
// Clear any default values
vals.clear();
// Recover from multi-value string
for (int i = 0; i < keys.size(); ++i)
{
if (keys[i] == string(arg_name))
{
string val_string(values[i]);
istringstream str_stream(val_string);
string::size_type old_pos = 0;
string::size_type new_pos = 0;
// Iterate <sep>-delimited values
value_t val;
while ((new_pos = val_string.find(sep, old_pos)) != string::npos)
{
if (new_pos != old_pos)
{
str_stream.width(new_pos - old_pos);
str_stream >> val;
vals.push_back(val);
}
// skip over delimiter
str_stream.ignore(1);
old_pos = new_pos + 1;
}
// Read last value
str_stream >> val;
vals.push_back(val);
}
}
}
}
/**
* The number of pairs parsed
*/
int parsed_argc()
{
return (int) keys.size();
}
/**
* Initialize device
*/
cudaError_t device_init()
{
cudaError_t error = cudaSuccess;
do
{
int deviceCount;
if (CUDA_PERROR(error = cudaGetDeviceCount(&deviceCount))) break;
if (deviceCount == 0) {
fprintf(stderr, "No devices supporting CUDA.\n");
exit(1);
}
if (device_id < 0)
{
get_cmd_line_argument("device", device_id);
}
if ((device_id > deviceCount - 1) || (device_id < 0))
{
device_id = 0;
}
if (CUDA_PERROR(error = cudaSetDevice(device_id))) break;
if (CUDA_PERROR(error = cudaMemGetInfo(&device_free_physmem, &device_total_physmem))) break;
if (CUDA_PERROR(error = cudaGetDeviceProperties(&device_prop, device_id))) break;
if (device_prop.major < 1) {
fprintf(stderr, "Device does not support CUDA.\n");
exit(1);
}
device_giga_bandwidth = float(device_prop.memoryBusWidth) * device_prop.memoryClockRate * 2 / 8 / 1000 / 1000;
} while (0);
return error;
}
//-------------------------------------------------------------------------
// Utility functions
//-------------------------------------------------------------------------
/// Tokenizes a comma-delimited list of string pairs delimited by ':'
static void tokenize(
std::vector<std::pair<std::string, std::string> > &tokens,
std::string const &str,
char delim = ',',
char sep = ':')
{
// Home-built to avoid Boost dependency
size_t s_idx = 0;
size_t d_idx = std::string::npos;
while (s_idx < str.size())
{
d_idx = str.find_first_of(delim, s_idx);
size_t end_idx = (d_idx != std::string::npos ? d_idx : str.size());
size_t sep_idx = str.find_first_of(sep, s_idx);
size_t offset = 1;
if (sep_idx == std::string::npos || sep_idx >= end_idx)
{
sep_idx = end_idx;
offset = 0;
}
std::pair<std::string, std::string> item(
str.substr(s_idx, sep_idx - s_idx),
str.substr(sep_idx + offset, end_idx - sep_idx - offset));
tokens.push_back(item);
s_idx = end_idx + 1;
}
}
/// Tokenizes a comma-delimited list of string pairs delimited by ':'
static void tokenize(
std::vector<std::string > &tokens,
std::string const &str,
char delim = ',',
char sep = ':')
{
std::vector<std::pair<std::string, std::string> > token_pairs;
tokenize(token_pairs, str, delim, sep);
for (auto const &tok : token_pairs)
{
tokens.push_back(tok.first);
}
}
};
} // namespace cutlass

View File

@ -1,83 +0,0 @@
/******************************************************************************
* Copyright (c) 2011-2017, NVIDIA CORPORATION. All rights reserved.
*
* Redistribution and use in source and binary forms, with or without
* modification, are not permitted.
*
* THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
* ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
* WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
* DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE FOR ANY
* DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
* (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
* LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
* ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
* (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
* SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
*
******************************************************************************/
#pragma once
/**
* \file
* \brief C++ exception semantics for CUDA error codes
*/
#include <iosfwd>
#include <cuda_runtime.h>
namespace cutlass {
/// C++ exception wrapper for CUDA \p cudaError_t
class cuda_exception : public std::exception
{
public:
/// Constructor
cuda_exception(
const char *msg = "",
cudaError_t err = cudaErrorUnknown)
:
msg(msg), err(err)
{}
/// Returns the explanatory string
const char *what() const noexcept
{
return msg;
}
/// Returns the underlying CUDA \p cudaError_t
cudaError_t cudaError() const
{
return err;
}
protected:
/// Explanatory string
const char *msg;
/// Underlying CUDA \p cudaError_t
cudaError_t err;
};
/// Writes a cudaError_t to an output stream
inline std::ostream & operator<<(std::ostream &out, cudaError_t result)
{
return out << cudaGetErrorString(result);
}
/// Writes a cuda_exception instance to an output stream
inline std::ostream & operator<<(std::ostream &out, cuda_exception const &e)
{
return out << e.what() << ": " << e.cudaError();
}
} // namespace cutlass

View File

@ -1,224 +0,0 @@
/******************************************************************************
* Copyright (c) 2011-2017, NVIDIA CORPORATION. All rights reserved.
*
* Redistribution and use in source and binary forms, with or without
* modification, are not permitted.
*
* THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
* ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
* WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
* DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE FOR ANY
* DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
* (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
* LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
* ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
* (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
* SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
*
******************************************************************************/
#pragma once
/**
* \file
* Utilities for interacting with the opaque CUDA __half type
*/
#include <stdint.h>
#include <cuda_fp16.h>
#include <iosfwd>
namespace cutlass {
/******************************************************************************
* half_t
******************************************************************************/
/**
* Host-based fp16 data type compatible and convertible with __half
*/
struct half_t
{
uint16_t __x;
/// Constructor from __half
half_t(const __half &other)
{
__x = reinterpret_cast<const uint16_t&>(other);
}
/// Constructor from integer
half_t(int a)
{
*this = half_t(float(a));
}
/// Constructor from float
half_t(float a)
{
uint32_t ia = *reinterpret_cast<uint32_t*>(&a);
uint16_t ir;
ir = (ia >> 16) & 0x8000;
if ((ia & 0x7f800000) == 0x7f800000)
{
if ((ia & 0x7fffffff) == 0x7f800000)
{
ir |= 0x7c00; /* infinity */
}
else
{
ir = 0x7fff; /* canonical NaN */
}
}
else if ((ia & 0x7f800000) >= 0x33000000)
{
int32_t shift = (int32_t) ((ia >> 23) & 0xff) - 127;
if (shift > 15)
{
ir |= 0x7c00; /* infinity */
}
else
{
ia = (ia & 0x007fffff) | 0x00800000; /* extract mantissa */
if (shift < -14)
{ /* denormal */
ir |= ia >> (-1 - shift);
ia = ia << (32 - (-1 - shift));
}
else
{ /* normal */
ir |= ia >> (24 - 11);
ia = ia << (32 - (24 - 11));
ir = ir + ((14 + shift) << 10);
}
/* IEEE-754 round to nearest of even */
if ((ia > 0x80000000) || ((ia == 0x80000000) && (ir & 1)))
{
ir++;
}
}
}
this->__x = ir;
}
/// Cast to __half
operator __half() const
{
return reinterpret_cast<const __half&>(__x);
}
/// Cast to float
operator float() const
{
int sign = ((this->__x >> 15) & 1);
int exp = ((this->__x >> 10) & 0x1f);
int mantissa = (this->__x & 0x3ff);
uint32_t f = 0;
if (exp > 0 && exp < 31)
{
// normal
exp += 112;
f = (sign << 31) | (exp << 23) | (mantissa << 13);
}
else if (exp == 0)
{
if (mantissa)
{
// subnormal
exp += 113;
while ((mantissa & (1 << 10)) == 0)
{
mantissa <<= 1;
exp--;
}
mantissa &= 0x3ff;
f = (sign << 31) | (exp << 23) | (mantissa << 13);
}
else
{
// zero
f = 0;
}
}
else if (exp == 31)
{
if (mantissa)
{
f = 0x7fffffff; // not a number
}
else
{
f = (0xff << 23) | (sign << 31); // inf
}
}
return *reinterpret_cast<float const *>(&f);
}
/// Get raw storage
uint16_t raw()
{
return this->__x;
}
/// Assignment by sum
bool operator ==(const half_t &other)
{
return (this->__x == other.__x);
}
/// Increment
half_t& operator +=(const half_t &rhs)
{
*this = half_t(float(*this) + float(rhs));
return *this;
}
/// Decrement
half_t& operator -=(const half_t &rhs)
{
*this = half_t(float(*this) - float(rhs));
return *this;
}
/// Multiply
half_t operator*(const half_t &other)
{
return half_t(float(*this) * float(other));
}
/// Multiply
half_t operator+(const half_t &other)
{
return half_t(float(*this) + float(other));
}
};
/******************************************************************************
* I/O stream overloads
******************************************************************************/
/// Insert formatted \p half_t into the output stream
std::ostream& operator<<(std::ostream &out, const half_t &x)
{
out << (float)x;
return out;
}
/// Insert formatted \p __half into the output stream
std::ostream& operator<<(std::ostream &out, const __half &x)
{
return out << half_t(x);
}
} // namespace cutlass

View File

@ -1,495 +0,0 @@
/******************************************************************************
* Copyright (c) 2011-2017, NVIDIA CORPORATION. All rights reserved.
*
* Redistribution and use in source and binary forms, with or without
* modification, are not permitted.
*
* THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
* ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
* WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
* DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE FOR ANY
* DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
* (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
* LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
* ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
* (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
* SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
*
******************************************************************************/
#pragma once
/**
* \file
* Matrix data structure providing basic CPU-based algorithms and
* operations that can be cloned and synchronized in GPU device memory
*/
#include <vector>
#include <fstream>
#include <cutlass/util/debug.h>
#include "../cutlass/util/matrix_transform.h"
#include "half.h"
namespace cutlass {
/**
* \brief Matrix data structure providing basic CPU-based algorithms and
* operations that be synchronized with a GPU-based replica
*/
template <typename value_t>
struct matrix
{
// Host value type (must be convertible to/from value_t)
typedef typename nv_std::conditional<
(nv_std::is_same<value_t, __half>::value), // If (value_t == __half) ...
half_t, // ... use half_t internally for host storage, else...
value_t>::type // ... use value_t directly
host_value_t;
//-----------------------------------------------------------------------------
// Data members
//-----------------------------------------------------------------------------
private:
/// M dimension (height in rows)
int _m;
/// N dimension (width in columns)
int _n;
/// Data array on host
std::vector<host_value_t> _h_data;
/// Clone of data array on GPU device
value_t *_d_data;
/// GPU Device identifier that clone synchronizes with
int _device_id;
public:
//-----------------------------------------------------------------------------
// Lifetime and synchronization
//-----------------------------------------------------------------------------
/**
* Constructor: zero-initializes the matrix.
*/
matrix(
int m, ///< Height of the matrix in rows
int n) ///< Width of the matrix in columns
:
_m(m),
_n(n),
_d_data(NULL),
_device_id(0)
{
_h_data.resize(_m * _n, 0);
CUDA_PERROR_EXIT(cudaMalloc((void ** )&_d_data, sizeof(value_t) * _m * _n));
CUDA_PERROR_EXIT(cudaGetDevice(&_device_id));
}
/// Destructor
~matrix()
{
if (_d_data)
{
CUDA_PERROR_EXIT(cudaFree(_d_data));
}
}
/**
* Synchronize the GPU-based replica with the current host-based matrix data
*/
void sync_device()
{
size_t bytes = _m * _n * sizeof(value_t);
CUDA_PERROR_EXIT(cudaMemcpy(_d_data, &_h_data[0], bytes, cudaMemcpyHostToDevice));
}
/**
* Synchronize the host-based replica with the current GPU-based matrix data
*/
void sync_host()
{
size_t bytes = _m * _n * sizeof(value_t);
CUDA_PERROR_EXIT(cudaMemcpy(&_h_data[0], _d_data, bytes, cudaMemcpyDeviceToHost));
}
//-----------------------------------------------------------------------------
// Inspectors
//-----------------------------------------------------------------------------
/**
* Return the height of the matrix, subject to the optional \p transpose_op
*/
int height(matrix_transform_t transpose_op = matrix_transform_t::NonTranspose) const
{
switch (transpose_op)
{
case matrix_transform_t::NonTranspose : return _m;
case matrix_transform_t::Transpose : return _n;
default: return -1;
}
}
/**
* Return the width of the matrix, subject to the optional \p transpose_op
*/
int width(matrix_transform_t transpose_op = matrix_transform_t::NonTranspose) const
{
switch (transpose_op)
{
case matrix_transform_t::NonTranspose : return _n;
case matrix_transform_t::Transpose : return _m;
default: return -1;
}
}
/**
* Return item at (x, y) coordinate of matrix, subject to the optional \p transform op
*/
host_value_t get(
int x,
int y,
matrix_transform_t transpose_op = matrix_transform_t::NonTranspose) const
{
switch (transpose_op)
{
case matrix_transform_t::NonTranspose : return _h_data[y + (x * _m)];
case matrix_transform_t::Transpose : return _h_data[x + (y * _m)];
default: return 0;
}
}
/**
* Return the distance (in items) within memory between elements of two
* consecutive columns which have the same row index, subject to the optional \p transform op
*/
int leading_dim(matrix_transform_t transpose_op = matrix_transform_t::NonTranspose) const
{
switch (transpose_op)
{
case matrix_transform_t::NonTranspose : return _m;
case matrix_transform_t::Transpose : return _n;
default: return 0;
}
}
/**
* Get host data pointer
*/
value_t* h_data()
{
return _h_data.data();
}
/**
* Get host data pointer
*/
value_t const* h_data() const
{
return _h_data.data();
}
/**
* Get device data pointer
*/
value_t const* d_data() const
{
return _d_data;
}
/**
* Get device data pointer
*/
value_t * d_data()
{
return _d_data;
}
//-----------------------------------------------------------------------------
// Initialization
//-----------------------------------------------------------------------------
/**
* Initialize matrix values with a 2D "ramp" defined as
* <tt>values(x, y) = (y * rs) + (x * cs)</tt>
*/
void fill_ramp(
host_value_t rs,
host_value_t cs)
{
for (int x = 0; x < _n; x++)
{
for (int y = 0; y < _m; y++)
{
_h_data[y + (x * _m)] = host_value_t((y * rs) + (x * cs));
}
}
}
/**
* Initialize matrix values such that all the elements of the principal diagonal
* are ones and all other elements are zeros
*/
void fill_identity()
{
for (int j = 0; j < _n; j++)
{
for (int i = 0; i < _m; i++)
{
_h_data[i + j * _m] = host_value_t(i == j ? 1 : 0);
}
}
}
/**
* Initialize matrix values using the random number \p generator. The
* \p generator reference is assumed to be a nullary functor that returns
* values convertible to the matrix \p value_t.
*/
template <typename T>
void fill_random(T & generator)
{
for (int j = 0; j < _n; j++)
{
for (int i = 0; i < _m; i++)
{
_h_data[i + j * _m] = (value_t) generator();
}
}
}
/**
* Element-wise matrix addition
*/
matrix & operator+=(matrix const &mat)
{
for (int j = 0; j < _n; j++)
{
for (int i = 0; i < _m; i++)
{
_h_data[i + j * _m] += mat._h_data[i + j * _m];
}
}
return *this;
}
/**
* Element-wise matrix subtraction
*/
matrix & operator-=(matrix const &mat)
{
for (int j = 0; j < _n; j++)
{
for (int i = 0; i < _m; i++)
{
_h_data[i + j * _m] -= mat._h_data[i + j * _m];
}
}
return *this;
}
//-----------------------------------------------------------------------------
// Output
//-----------------------------------------------------------------------------
/**
* Prints matrix in CSV to output stream
*/
template <typename _hv_t>
std::ostream & write_matrix(std::ostream &out, _hv_t)
{
for (int i = 0; i < _m; i++)
{
for (int j = 0; j < _n; j++)
{
out << (j ? "," : "") << _h_data[i + j * _m];
}
out << "\n";
}
return out;
}
/**
* Prints matrix in CSV to output stream
*/
std::ostream & write_matrix(std::ostream &out, int8_t)
{
for (int i = 0; i < _m; i++)
{
for (int j = 0; j < _n; j++)
{
out << (j ? "," : "") << int32_t(_h_data[i + j * _m]);
}
out << "\n";
}
return out;
}
/**
* Prints matrix in CSV to output stream
*/
std::ostream & write_matrix(std::ostream &out)
{
return write_matrix(out, _h_data[0]);
}
//-----------------------------------------------------------------------------
// Floating point "almost-equal" utilities
//-----------------------------------------------------------------------------
static bool almost_equal_ulps(half_t a, half_t b, int max_ulps)
{
if (a == b)
return true;
int32_t int_diff = abs(a.raw() - b.raw());
if (int_diff <= max_ulps)
return true;
return false;
}
static bool almost_equal_ulps(float a, float b, int max_ulps)
{
if (a == b)
return true;
int32_t int_diff = abs(*(int32_t*)&a - *(int32_t*)&b);
if (int_diff <= max_ulps)
return true;
return false;
}
static bool almost_equal_ulps(double a, double b, int max_ulps)
{
if (a == b)
return true;
int64_t int_diff = abs(*(int64_t*)&a - *(int64_t*)&b);
if (int_diff <= max_ulps)
return true;
return false;
}
static bool almost_equal_ulps(int32_t a, int32_t b, int max_ulps)
{
return (a == b);
}
//-----------------------------------------------------------------------------
// matrix operations
//-----------------------------------------------------------------------------
/**
* Returns matrix equality
*/
bool operator==(const matrix<value_t> &mat) const
{
int max_ulps = 30;
if (_m != mat._m || _n != mat._n)
{
fprintf(stderr, "Error: dimension mismatch during matrix comparison.\n"); exit(1);
}
for (int j = 0; j < _n; j++)
{
for (int i = 0; i < _m; i++)
{
if (!almost_equal_ulps(_h_data[i + j * _m], mat._h_data[i + j * _m], max_ulps))
{
return false;
}
}
}
return true;
}
/**
* Returns matrix inequality
*/
bool operator!=(const matrix<value_t> &mat) const
{
return !(*this == mat);
}
/**
* Computes this = (alpha * op(A) * op(B)) + (beta * this), specialized for gemm_nn
*/
template <typename multiplicand_t>
void gemm(
matrix_transform_t transform_a,
matrix_transform_t transform_b,
host_value_t alpha,
const matrix<multiplicand_t> &A,
const matrix<multiplicand_t> &B,
host_value_t beta)
{
// Sanity check dimensions
if ((_m != A.height(transform_a)) ||
(_n != B.width(transform_b)) ||
(A.width(transform_a) != B.height(transform_b)))
{
fprintf(stderr, "Error: dimension mismatch during gemm.\n");
exit(1);
}
int M = A.height(transform_a);
int K = A.width(transform_a);
int N = B.width(transform_b);
// Even the host-side implementation utilizes a blocking structure to improve
// verification performance
int DimBlockM = (M % 16 == 0) ? 16 : 1;
int DimBlockN = (N % 16 == 0) ? 16 : 1;
for (int i = 0; i < M; i += DimBlockM)
{
for (int j = 0; j < N; j += DimBlockN)
{
for (int block_y = 0; block_y < DimBlockM; block_y++)
{
for (int block_x = 0; block_x < DimBlockN; block_x++)
{
int y = i + block_y;
int x = j + block_x;
host_value_t accum(0);
for (int k = 0; k < K; k++)
{
accum += host_value_t(A.get(k, y, transform_a)) * host_value_t(B.get(x, k, transform_b));
}
_h_data[y + x * M] = (alpha * accum) + (beta * _h_data[y + x * M]);
}
}
}
}
}
};
} // namespace cutlass

View File

@ -1,99 +0,0 @@
/******************************************************************************
* Copyright (c) 2011-2017, NVIDIA CORPORATION. All rights reserved.
*
* Redistribution and use in source and binary forms, with or without
* modification, are not permitted.
*
* THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
* ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
* WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
* DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE FOR ANY
* DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
* (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
* LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
* ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
* (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
* SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
*
******************************************************************************/
#pragma once
/**
* \file
* GPU kernel timer
*/
#include <cuda_runtime.h>
#include <cutlass/util/debug.h>
namespace cutlass {
/******************************************************************************
* gpu_timer
******************************************************************************/
/**
* GPU event-based timer
*/
struct gpu_timer
{
cudaEvent_t _start;
cudaEvent_t _stop;
gpu_timer()
{
CUDA_PERROR_EXIT(cudaEventCreate(&_start));
CUDA_PERROR_EXIT(cudaEventCreate(&_stop));
}
~gpu_timer()
{
CUDA_PERROR_EXIT(cudaEventDestroy(_start));
CUDA_PERROR_EXIT(cudaEventDestroy(_stop));
}
void start()
{
CUDA_PERROR_EXIT(cudaEventRecord(_start, 0));
}
void stop()
{
CUDA_PERROR_EXIT(cudaEventRecord(_stop, 0));
}
float elapsed_millis()
{
float elapsed = 0.0;
CUDA_PERROR_EXIT(cudaEventSynchronize(_stop));
CUDA_PERROR_EXIT(cudaEventElapsedTime(&elapsed, _start, _stop));
return elapsed;
}
};
/******************************************************************************
* sleep_millis
******************************************************************************/
#ifdef _WIN32
#include <windows.h>
void sleep_millis(unsigned milliseconds)
{
Sleep(milliseconds);
}
#else
#include <unistd.h>
void sleep_millis(unsigned milliseconds)
{
usleep(milliseconds * 1000); // takes microseconds
}
#endif
} // namespace cutlass

View File

@ -1,155 +0,0 @@
/******************************************************************************
* Copyright (c) 2011-2017, NVIDIA CORPORATION. All rights reserved.
*
* Redistribution and use in source and binary forms, with or without
* modification, are not permitted.
*
* THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
* ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
* WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
* DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE FOR ANY
* DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
* (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
* LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
* ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
* (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
* SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
*
******************************************************************************/
#pragma once
/**
* \file
* \brief Utilities for converting between types and assessing traits
*/
#include "half.h"
namespace cutlass {
/******************************************************************************
* Float conversion utilities
******************************************************************************/
/// Convert float to value type
template <typename value_t>
value_t from_float(float val)
{
return value_t(val);
}
/// Convert float to value type (__half specialization)
template <>
__half from_float<__half>(float val)
{
return half_t(val);
}
/******************************************************************************
* Type conversion utilities
******************************************************************************/
/// Member \p type is defined as the signed integer type having the same size as \p T
template <typename T>
struct integer_alias;
template <>
struct integer_alias<int8_t> {
using type = int8_t;
};
template <>
struct integer_alias<half_t> {
using type = int16_t;
};
template <>
struct integer_alias<__half> {
using type = int16_t;
};
template <>
struct integer_alias<float> {
using type = int32_t;
};
template <>
struct integer_alias<int> {
using type = int32_t;
};
template <>
struct integer_alias<double> {
using type = int64_t;
};
/******************************************************************************
* Type-info utilities
******************************************************************************/
/// Returns a string to prefix 'gemm' to construct CUBLAS-like kernel names
template <math_operation_class_t math_op, typename value_t, typename accum_t> char const *to_prefix_string();
template <> char const *to_prefix_string<math_operation_class_t::scalar, half_t, half_t>() {
return "H";
}
template <> char const *to_prefix_string<math_operation_class_t::scalar, __half, __half>() {
return "H";
}
template <> char const *to_prefix_string<math_operation_class_t::scalar, float, float>() {
return "S";
}
template <> char const *to_prefix_string<math_operation_class_t::matrix, __half, __half>() {
return "WmmaH";
}
template <> char const *to_prefix_string<math_operation_class_t::matrix, __half, float>() {
return "WmmaS";
}
template <> char const *to_prefix_string<math_operation_class_t::scalar, double, double>() {
return "D";
}
template <> char const *to_prefix_string<math_operation_class_t::scalar, int8_t, int32_t>() {
return "I";
}
/******************************************************************************
* Maps value_t to the minimum vector size used to load operand
******************************************************************************/
template <typename T>
struct operand_load_type;
template <>
struct operand_load_type<int8_t> { using type = int32_t; };
template <typename T>
struct operand_load_type { using type = T; };
/******************************************************************************
* Minimum alignment requirement, if any, determined from value_t.
******************************************************************************/
template <typename value_t>
struct gemm_alignment_requirement;
template <>
struct gemm_alignment_requirement<uint8_t> { static const int value = 4; };
template <typename value_t>
struct gemm_alignment_requirement { static const int value = 0; };
} // namespace cutlass

1
docs/_config.yml Normal file
View File

@ -0,0 +1 @@
theme: jekyll-theme-minimal

View File

@ -0,0 +1,145 @@
<!DOCTYPE html PUBLIC "-//W3C//DTD XHTML 1.0 Transitional//EN" "http://www.w3.org/TR/xhtml1/DTD/xhtml1-transitional.dtd">
<html xmlns="http://www.w3.org/1999/xhtml">
<head>
<meta http-equiv="Content-Type" content="text/xhtml;charset=UTF-8"/>
<meta http-equiv="X-UA-Compatible" content="IE=9"/>
<meta name="generator" content="Doxygen 1.8.11"/>
<title>CUTLASS: aligned_buffer.h File Reference</title>
<link href="tabs.css" rel="stylesheet" type="text/css"/>
<script type="text/javascript" src="jquery.js"></script>
<script type="text/javascript" src="dynsections.js"></script>
<link href="search/search.css" rel="stylesheet" type="text/css"/>
<script type="text/javascript" src="search/searchdata.js"></script>
<script type="text/javascript" src="search/search.js"></script>
<script type="text/javascript">
$(document).ready(function() { init_search(); });
</script>
<script type="text/x-mathjax-config">
MathJax.Hub.Config({
extensions: ["tex2jax.js"],
jax: ["input/TeX","output/HTML-CSS"],
});
</script><script type="text/javascript" src="http://cdn.mathjax.org/mathjax/latest/MathJax.js"></script>
<link href="doxygen.css" rel="stylesheet" type="text/css" />
</head>
<body>
<div id="top"><!-- do not remove this div, it is closed by doxygen! -->
<div id="titlearea">
<table cellspacing="0" cellpadding="0">
<tbody>
<tr style="height: 56px;">
<td id="projectlogo"><img alt="Logo" src="cutlass-logo-small.png"/></td>
<td id="projectalign" style="padding-left: 0.5em;">
<div id="projectname">CUTLASS
</div>
<div id="projectbrief">CUDA Templates for Linear Algebra Subroutines and Solvers</div>
</td>
</tr>
</tbody>
</table>
</div>
<!-- end header part -->
<!-- Generated by Doxygen 1.8.11 -->
<script type="text/javascript">
var searchBox = new SearchBox("searchBox", "search",false,'Search');
</script>
<div id="navrow1" class="tabs">
<ul class="tablist">
<li><a href="index.html"><span>Main&#160;Page</span></a></li>
<li><a href="modules.html"><span>Modules</span></a></li>
<li><a href="namespaces.html"><span>Namespaces</span></a></li>
<li><a href="annotated.html"><span>Classes</span></a></li>
<li class="current"><a href="files.html"><span>Files</span></a></li>
<li>
<div id="MSearchBox" class="MSearchBoxInactive">
<span class="left">
<img id="MSearchSelect" src="search/mag_sel.png"
onmouseover="return searchBox.OnSearchSelectShow()"
onmouseout="return searchBox.OnSearchSelectHide()"
alt=""/>
<input type="text" id="MSearchField" value="Search" accesskey="S"
onfocus="searchBox.OnSearchFieldFocus(true)"
onblur="searchBox.OnSearchFieldFocus(false)"
onkeyup="searchBox.OnSearchFieldChange(event)"/>
</span><span class="right">
<a id="MSearchClose" href="javascript:searchBox.CloseResultsWindow()"><img id="MSearchCloseImg" border="0" src="search/close.png" alt=""/></a>
</span>
</div>
</li>
</ul>
</div>
<div id="navrow2" class="tabs2">
<ul class="tablist">
<li><a href="files.html"><span>File&#160;List</span></a></li>
<li><a href="globals.html"><span>File&#160;Members</span></a></li>
</ul>
</div>
<!-- window showing the filter options -->
<div id="MSearchSelectWindow"
onmouseover="return searchBox.OnSearchSelectShow()"
onmouseout="return searchBox.OnSearchSelectHide()"
onkeydown="return searchBox.OnSearchSelectKey(event)">
</div>
<!-- iframe showing the search results (closed by default) -->
<div id="MSearchResultsWindow">
<iframe src="javascript:void(0)" frameborder="0"
name="MSearchResults" id="MSearchResults">
</iframe>
</div>
<div id="nav-path" class="navpath">
<ul>
<li class="navelem"><a class="el" href="dir_d44c64559bbebec7f509842c48db8b23.html">include</a></li><li class="navelem"><a class="el" href="dir_6baf2bb612a2f0daa69af3101ede80a1.html">cutlass</a></li> </ul>
</div>
</div><!-- top -->
<div class="header">
<div class="summary">
<a href="#nested-classes">Classes</a> &#124;
<a href="#namespaces">Namespaces</a> </div>
<div class="headertitle">
<div class="title">aligned_buffer.h File Reference</div> </div>
</div><!--header-->
<div class="contents">
<p>AlignedBuffer is a container for trivially copyable elements suitable for use in unions and shared memory.
<a href="#details">More...</a></p>
<div class="textblock"><code>#include &quot;<a class="el" href="cutlass_8h_source.html">cutlass/cutlass.h</a>&quot;</code><br />
<code>#include &quot;<a class="el" href="array_8h_source.html">cutlass/array.h</a>&quot;</code><br />
</div><div class="textblock"><div class="dynheader">
Include dependency graph for aligned_buffer.h:</div>
<div class="dyncontent">
<div class="center"><img src="aligned__buffer_8h__incl.png" border="0" usemap="#aligned__buffer_8h" alt=""/></div>
<map name="aligned__buffer_8h" id="aligned__buffer_8h">
</map>
</div>
</div><div class="textblock"><div class="dynheader">
This graph shows which files directly or indirectly include this file:</div>
<div class="dyncontent">
<div class="center"><img src="aligned__buffer_8h__dep__incl.png" border="0" usemap="#aligned__buffer_8hdep" alt=""/></div>
<map name="aligned__buffer_8hdep" id="aligned__buffer_8hdep">
</map>
</div>
</div>
<p><a href="aligned__buffer_8h_source.html">Go to the source code of this file.</a></p>
<table class="memberdecls">
<tr class="heading"><td colspan="2"><h2 class="groupheader"><a name="nested-classes"></a>
Classes</h2></td></tr>
<tr class="memitem:"><td class="memItemLeft" align="right" valign="top">struct &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="structcutlass_1_1AlignedBuffer.html">cutlass::AlignedBuffer&lt; T, N, Align &gt;</a></td></tr>
<tr class="memdesc:"><td class="mdescLeft">&#160;</td><td class="mdescRight">Modifies semantics of cutlass::Array&lt;&gt; to provide guaranteed alignment. <a href="structcutlass_1_1AlignedBuffer.html#details">More...</a><br /></td></tr>
<tr class="separator:"><td class="memSeparator" colspan="2">&#160;</td></tr>
</table><table class="memberdecls">
<tr class="heading"><td colspan="2"><h2 class="groupheader"><a name="namespaces"></a>
Namespaces</h2></td></tr>
<tr class="memitem:namespacecutlass"><td class="memItemLeft" align="right" valign="top"> &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="namespacecutlass.html">cutlass</a></td></tr>
<tr class="separator:"><td class="memSeparator" colspan="2">&#160;</td></tr>
</table>
</div><!-- contents -->
<!-- start footer part -->
<hr class="footer"/><address class="footer"><small>
Generated by &#160;<a href="http://www.doxygen.org/index.html">
<img class="footer" src="doxygen.png" alt="doxygen"/>
</a> 1.8.11
</small></address>
</body>
</html>

View File

@ -0,0 +1 @@
6cbc6b81ede44b5f08afd4f4519d56d1

View File

@ -0,0 +1 @@
b26c62930ff7668b89f2ee6624e0be3a

File diff suppressed because one or more lines are too long

867
docs/annotated.html Normal file
View File

@ -0,0 +1,867 @@
<!DOCTYPE html PUBLIC "-//W3C//DTD XHTML 1.0 Transitional//EN" "http://www.w3.org/TR/xhtml1/DTD/xhtml1-transitional.dtd">
<html xmlns="http://www.w3.org/1999/xhtml">
<head>
<meta http-equiv="Content-Type" content="text/xhtml;charset=UTF-8"/>
<meta http-equiv="X-UA-Compatible" content="IE=9"/>
<meta name="generator" content="Doxygen 1.8.11"/>
<title>CUTLASS: Class List</title>
<link href="tabs.css" rel="stylesheet" type="text/css"/>
<script type="text/javascript" src="jquery.js"></script>
<script type="text/javascript" src="dynsections.js"></script>
<link href="search/search.css" rel="stylesheet" type="text/css"/>
<script type="text/javascript" src="search/searchdata.js"></script>
<script type="text/javascript" src="search/search.js"></script>
<script type="text/javascript">
$(document).ready(function() { init_search(); });
</script>
<script type="text/x-mathjax-config">
MathJax.Hub.Config({
extensions: ["tex2jax.js"],
jax: ["input/TeX","output/HTML-CSS"],
});
</script><script type="text/javascript" src="http://cdn.mathjax.org/mathjax/latest/MathJax.js"></script>
<link href="doxygen.css" rel="stylesheet" type="text/css" />
</head>
<body>
<div id="top"><!-- do not remove this div, it is closed by doxygen! -->
<div id="titlearea">
<table cellspacing="0" cellpadding="0">
<tbody>
<tr style="height: 56px;">
<td id="projectlogo"><img alt="Logo" src="cutlass-logo-small.png"/></td>
<td id="projectalign" style="padding-left: 0.5em;">
<div id="projectname">CUTLASS
</div>
<div id="projectbrief">CUDA Templates for Linear Algebra Subroutines and Solvers</div>
</td>
</tr>
</tbody>
</table>
</div>
<!-- end header part -->
<!-- Generated by Doxygen 1.8.11 -->
<script type="text/javascript">
var searchBox = new SearchBox("searchBox", "search",false,'Search');
</script>
<div id="navrow1" class="tabs">
<ul class="tablist">
<li><a href="index.html"><span>Main&#160;Page</span></a></li>
<li><a href="modules.html"><span>Modules</span></a></li>
<li><a href="namespaces.html"><span>Namespaces</span></a></li>
<li class="current"><a href="annotated.html"><span>Classes</span></a></li>
<li><a href="files.html"><span>Files</span></a></li>
<li>
<div id="MSearchBox" class="MSearchBoxInactive">
<span class="left">
<img id="MSearchSelect" src="search/mag_sel.png"
onmouseover="return searchBox.OnSearchSelectShow()"
onmouseout="return searchBox.OnSearchSelectHide()"
alt=""/>
<input type="text" id="MSearchField" value="Search" accesskey="S"
onfocus="searchBox.OnSearchFieldFocus(true)"
onblur="searchBox.OnSearchFieldFocus(false)"
onkeyup="searchBox.OnSearchFieldChange(event)"/>
</span><span class="right">
<a id="MSearchClose" href="javascript:searchBox.CloseResultsWindow()"><img id="MSearchCloseImg" border="0" src="search/close.png" alt=""/></a>
</span>
</div>
</li>
</ul>
</div>
<div id="navrow2" class="tabs2">
<ul class="tablist">
<li class="current"><a href="annotated.html"><span>Class&#160;List</span></a></li>
<li><a href="classes.html"><span>Class&#160;Index</span></a></li>
<li><a href="inherits.html"><span>Class&#160;Hierarchy</span></a></li>
<li><a href="functions.html"><span>Class&#160;Members</span></a></li>
</ul>
</div>
</div><!-- top -->
<!-- window showing the filter options -->
<div id="MSearchSelectWindow"
onmouseover="return searchBox.OnSearchSelectShow()"
onmouseout="return searchBox.OnSearchSelectHide()"
onkeydown="return searchBox.OnSearchSelectKey(event)">
</div>
<!-- iframe showing the search results (closed by default) -->
<div id="MSearchResultsWindow">
<iframe src="javascript:void(0)" frameborder="0"
name="MSearchResults" id="MSearchResults">
</iframe>
</div>
<div class="header">
<div class="headertitle">
<div class="title">Class List</div> </div>
</div><!--header-->
<div class="contents">
<div class="textblock">Here are the classes, structs, unions and interfaces with brief descriptions:</div><div class="directory">
<div class="levels">[detail level <span onclick="javascript:toggleLevel(1);">1</span><span onclick="javascript:toggleLevel(2);">2</span><span onclick="javascript:toggleLevel(3);">3</span><span onclick="javascript:toggleLevel(4);">4</span><span onclick="javascript:toggleLevel(5);">5</span><span onclick="javascript:toggleLevel(6);">6</span>]</div><table class="directory">
<tr id="row_0_" class="even"><td class="entry"><span style="width:0px;display:inline-block;">&#160;</span><span id="arr_0_" class="arrow" onclick="toggleFolder('0_')">&#9658;</span><span class="icona"><span class="icon">N</span></span><a class="el" href="namespacecutlass.html" target="_self">cutlass</a></td><td class="desc"></td></tr>
<tr id="row_0_0_" style="display:none;"><td class="entry"><span style="width:16px;display:inline-block;">&#160;</span><span id="arr_0_0_" class="arrow" onclick="toggleFolder('0_0_')">&#9658;</span><span class="icona"><span class="icon">N</span></span><a class="el" href="namespacecutlass_1_1arch.html" target="_self">arch</a></td><td class="desc"></td></tr>
<tr id="row_0_0_0_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1arch_1_1Mma.html" target="_self">Mma</a></td><td class="desc">Matrix multiply-add operation </td></tr>
<tr id="row_0_0_1_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1arch_1_1Mma_3_01gemm_1_1GemmShape_3_011_00_011_00_011_01_4_00_011_00_01complex_30fa42e1ad201df010637cd22fc070a1.html" target="_self">Mma&lt; gemm::GemmShape&lt; 1, 1, 1 &gt;, 1, complex&lt; double &gt;, LayoutA, complex&lt; double &gt;, LayoutB, complex&lt; double &gt;, LayoutC, OpMultiplyAdd &gt;</a></td><td class="desc">Matrix multiply-add operation </td></tr>
<tr id="row_0_0_2_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1arch_1_1Mma_3_01gemm_1_1GemmShape_3_011_00_011_00_011_01_4_00_011_00_01complex_48b3a43bc03fff93a111ac01abe7e40d.html" target="_self">Mma&lt; gemm::GemmShape&lt; 1, 1, 1 &gt;, 1, complex&lt; double &gt;, LayoutA, double, LayoutB, complex&lt; double &gt;, LayoutC, OpMultiplyAdd &gt;</a></td><td class="desc">Matrix multiply-add operation </td></tr>
<tr id="row_0_0_3_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1arch_1_1Mma_3_01gemm_1_1GemmShape_3_011_00_011_00_011_01_4_00_011_00_01complex_76f9d24016e1b4167b16f4d7628c9546.html" target="_self">Mma&lt; gemm::GemmShape&lt; 1, 1, 1 &gt;, 1, complex&lt; float &gt;, LayoutA, complex&lt; float &gt;, LayoutB, complex&lt; float &gt;, LayoutC, OpMultiplyAdd &gt;</a></td><td class="desc">Matrix multiply-add operation </td></tr>
<tr id="row_0_0_4_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1arch_1_1Mma_3_01gemm_1_1GemmShape_3_011_00_011_00_011_01_4_00_011_00_01complex_f1c9d2ee842455cd0c5b71d56108d468.html" target="_self">Mma&lt; gemm::GemmShape&lt; 1, 1, 1 &gt;, 1, complex&lt; float &gt;, LayoutA, float, LayoutB, complex&lt; float &gt;, LayoutC, OpMultiplyAdd &gt;</a></td><td class="desc">Matrix multiply-add operation </td></tr>
<tr id="row_0_0_5_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1arch_1_1Mma_3_01gemm_1_1GemmShape_3_011_00_011_00_011_01_4_00_011_00_01double_070b94670e040ed5855e5b42d5ca8a443.html" target="_self">Mma&lt; gemm::GemmShape&lt; 1, 1, 1 &gt;, 1, double, LayoutA, complex&lt; double &gt;, LayoutB, complex&lt; double &gt;, LayoutC, OpMultiplyAdd &gt;</a></td><td class="desc">Matrix multiply-add operation </td></tr>
<tr id="row_0_0_6_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1arch_1_1Mma_3_01gemm_1_1GemmShape_3_011_00_011_00_011_01_4_00_011_00_01double_0aa57e6a2e6b5da37d10688bf99419a23.html" target="_self">Mma&lt; gemm::GemmShape&lt; 1, 1, 1 &gt;, 1, double, LayoutA, double, LayoutB, double, LayoutC, OpMultiplyAdd &gt;</a></td><td class="desc">Matrix multiply-add operation </td></tr>
<tr id="row_0_0_7_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1arch_1_1Mma_3_01gemm_1_1GemmShape_3_011_00_011_00_011_01_4_00_011_00_01ElementAb6e65b2cf5ede7f41cb070a767158dee.html" target="_self">Mma&lt; gemm::GemmShape&lt; 1, 1, 1 &gt;, 1, ElementA, LayoutA, ElementB, LayoutB, ElementC, LayoutC, Operator &gt;</a></td><td class="desc">Matrix multiply-add operation - specialized for 1x1x1x1 matrix multiply operation </td></tr>
<tr id="row_0_0_8_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1arch_1_1Mma_3_01gemm_1_1GemmShape_3_011_00_011_00_011_01_4_00_011_00_01float_00e3e12e263df6506b8cf06c3f4d478b8e.html" target="_self">Mma&lt; gemm::GemmShape&lt; 1, 1, 1 &gt;, 1, float, LayoutA, complex&lt; float &gt;, LayoutB, complex&lt; float &gt;, LayoutC, OpMultiplyAdd &gt;</a></td><td class="desc">Matrix multiply-add operation </td></tr>
<tr id="row_0_0_9_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1arch_1_1Mma_3_01gemm_1_1GemmShape_3_011_00_011_00_011_01_4_00_011_00_01float_004bb3fd76ca2af7b3210676fa9644d95b.html" target="_self">Mma&lt; gemm::GemmShape&lt; 1, 1, 1 &gt;, 1, float, LayoutA, float, LayoutB, float, LayoutC, OpMultiplyAdd &gt;</a></td><td class="desc">Matrix multiply-add operation </td></tr>
<tr id="row_0_0_10_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1arch_1_1Mma_3_01gemm_1_1GemmShape_3_011_00_011_00_011_01_4_00_011_00_01half__t_4f30ee91f7bb3844ff7579c68d078818.html" target="_self">Mma&lt; gemm::GemmShape&lt; 1, 1, 1 &gt;, 1, half_t, LayoutA, half_t, LayoutB, float, LayoutC, OpMultiplyAdd &gt;</a></td><td class="desc">Matrix multiply-add operation </td></tr>
<tr id="row_0_0_11_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1arch_1_1Mma_3_01gemm_1_1GemmShape_3_011_00_011_00_011_01_4_00_011_00_01int_00_00b2dff9ce8caad9aff5bc6a355539161.html" target="_self">Mma&lt; gemm::GemmShape&lt; 1, 1, 1 &gt;, 1, int, LayoutA, int, LayoutB, int, LayoutC, OpMultiplyAdd &gt;</a></td><td class="desc">Matrix multiply-add operation </td></tr>
<tr id="row_0_0_12_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1arch_1_1Mma_3_01gemm_1_1GemmShape_3_011_00_011_00_012_01_4_00_011_00_01int16__t8c4bac365710598317a69c489f7239db.html" target="_self">Mma&lt; gemm::GemmShape&lt; 1, 1, 2 &gt;, 1, int16_t, layout::RowMajor, int16_t, layout::ColumnMajor, int, LayoutC, OpMultiplyAdd &gt;</a></td><td class="desc">Matrix multiply-add operation </td></tr>
<tr id="row_0_0_13_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1arch_1_1Mma_3_01gemm_1_1GemmShape_3_011_00_011_00_014_01_4_00_011_00_01int8__t_a1ef6624fc8c10126f17f4ee88283d72.html" target="_self">Mma&lt; gemm::GemmShape&lt; 1, 1, 4 &gt;, 1, int8_t, LayoutA, int8_t, LayoutB, int, LayoutC, OpMultiplyAdd &gt;</a></td><td class="desc">Matrix multiply-add operation </td></tr>
<tr id="row_0_0_14_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1arch_1_1Mma_3_01gemm_1_1GemmShape_3_011_00_012_00_011_01_4_00_011_00_01half__t_f3dc2e59f857ada163d1e0781ea8f391.html" target="_self">Mma&lt; gemm::GemmShape&lt; 1, 2, 1 &gt;, 1, half_t, LayoutA, half_t, LayoutB, half_t, layout::RowMajor, OpMultiplyAdd &gt;</a></td><td class="desc">Matrix multiply-add operation </td></tr>
<tr id="row_0_0_15_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1arch_1_1Mma_3_01gemm_1_1GemmShape_3_0116_00_0116_00_014_01_4_00_0132_00_01half_0bcc4d05f9811035f08cc1b7f0154a4d.html" target="_self">Mma&lt; gemm::GemmShape&lt; 16, 16, 4 &gt;, 32, half_t, LayoutA, half_t, LayoutB, ElementC, LayoutC, Operator &gt;</a></td><td class="desc">Matrix multiply-add operation specialized for the entire warp </td></tr>
<tr id="row_0_0_16_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1arch_1_1Mma_3_01gemm_1_1GemmShape_3_0116_00_018_00_018_01_4_00_0132_00_01half__02a3f19a78995f97d793a668e0e4d4f0.html" target="_self">Mma&lt; gemm::GemmShape&lt; 16, 8, 8 &gt;, 32, half_t, layout::RowMajor, half_t, layout::ColumnMajor, float, layout::RowMajor, OpMultiplyAdd &gt;</a></td><td class="desc">Matrix multiply-add operation: F32 = F16 * F16 + F32 </td></tr>
<tr id="row_0_0_17_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1arch_1_1Mma_3_01gemm_1_1GemmShape_3_0116_00_018_00_018_01_4_00_0132_00_01half__96363097c47b056f0ca1911afd7f8b7a.html" target="_self">Mma&lt; gemm::GemmShape&lt; 16, 8, 8 &gt;, 32, half_t, layout::RowMajor, half_t, layout::ColumnMajor, half_t, layout::RowMajor, OpMultiplyAdd &gt;</a></td><td class="desc">Matrix multiply-add operation - F16 = F16 * F16 + F16 </td></tr>
<tr id="row_0_0_18_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1arch_1_1Mma_3_01gemm_1_1GemmShape_3_012_00_011_00_011_01_4_00_011_00_01half__t_8cf78649807b93684f3d431bfa34ee28.html" target="_self">Mma&lt; gemm::GemmShape&lt; 2, 1, 1 &gt;, 1, half_t, LayoutA, half_t, LayoutB, half_t, LayoutC, OpMultiplyAdd &gt;</a></td><td class="desc">Matrix multiply-add operation </td></tr>
<tr id="row_0_0_19_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1arch_1_1Mma_3_01gemm_1_1GemmShape_3_012_00_012_00_011_01_4_00_011_00_01half__t_ccde11d1bbbdab3702772ce44eb9729a.html" target="_self">Mma&lt; gemm::GemmShape&lt; 2, 2, 1 &gt;, 1, half_t, layout::ColumnMajor, half_t, layout::RowMajor, half_t, layout::ColumnMajor, OpMultiplyAdd &gt;</a></td><td class="desc">Matrix multiply-add operation </td></tr>
<tr id="row_0_0_20_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1arch_1_1Mma_3_01gemm_1_1GemmShape_3_012_00_012_00_011_01_4_00_011_00_01half__t_c07cc6439298fa5486a719e577be2538.html" target="_self">Mma&lt; gemm::GemmShape&lt; 2, 2, 1 &gt;, 1, half_t, layout::ColumnMajor, half_t, layout::RowMajor, half_t, layout::RowMajor, OpMultiplyAdd &gt;</a></td><td class="desc">Matrix multiply-add operation </td></tr>
<tr id="row_0_0_21_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1arch_1_1Mma_3_01gemm_1_1GemmShape_3_018_00_018_00_01128_01_4_00_0132_00_01uint15918972b95027764b3a849b03075ed2b.html" target="_self">Mma&lt; gemm::GemmShape&lt; 8, 8, 128 &gt;, 32, uint1b_t, layout::RowMajor, uint1b_t, layout::ColumnMajor, int, layout::RowMajor, OpXorPopc &gt;</a></td><td class="desc">Matrix multiply-add operation </td></tr>
<tr id="row_0_0_22_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1arch_1_1Mma_3_01gemm_1_1GemmShape_3_018_00_018_00_0116_01_4_00_0132_00_01int8__927179f46017ea5f58f859f1196c4829.html" target="_self">Mma&lt; gemm::GemmShape&lt; 8, 8, 16 &gt;, 32, int8_t, layout::RowMajor, int8_t, layout::ColumnMajor, int, layout::RowMajor, OpMultiplyAdd &gt;</a></td><td class="desc">Matrix multiply-add operation: S32 = S8 * S8 + S32 </td></tr>
<tr id="row_0_0_23_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1arch_1_1Mma_3_01gemm_1_1GemmShape_3_018_00_018_00_0116_01_4_00_0132_00_01int8__8ebae0cbdf333fddfe5c24d35ebe8e02.html" target="_self">Mma&lt; gemm::GemmShape&lt; 8, 8, 16 &gt;, 32, int8_t, layout::RowMajor, int8_t, layout::ColumnMajor, int, layout::RowMajor, OpMultiplyAddSaturate &gt;</a></td><td class="desc">Matrix multiply-add operation: S32 = S8 * S8 + S32 </td></tr>
<tr id="row_0_0_24_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1arch_1_1Mma_3_01gemm_1_1GemmShape_3_018_00_018_00_0116_01_4_00_0132_00_01int8__5299c9c90c8f2f521be0c8cec1c3eb08.html" target="_self">Mma&lt; gemm::GemmShape&lt; 8, 8, 16 &gt;, 32, int8_t, layout::RowMajor, uint8_t, layout::ColumnMajor, int, layout::RowMajor, OpMultiplyAdd &gt;</a></td><td class="desc">Matrix multiply-add operation: S32 = S8 * U8 + S32 </td></tr>
<tr id="row_0_0_25_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1arch_1_1Mma_3_01gemm_1_1GemmShape_3_018_00_018_00_0116_01_4_00_0132_00_01int8__f083347e265b1e9eea5572d86ddb6bf9.html" target="_self">Mma&lt; gemm::GemmShape&lt; 8, 8, 16 &gt;, 32, int8_t, layout::RowMajor, uint8_t, layout::ColumnMajor, int, layout::RowMajor, OpMultiplyAddSaturate &gt;</a></td><td class="desc">Matrix multiply-add operation: S32 = S8 * U8 + S32 </td></tr>
<tr id="row_0_0_26_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1arch_1_1Mma_3_01gemm_1_1GemmShape_3_018_00_018_00_0116_01_4_00_0132_00_01uint8_a62aa63a212985df306fb27e8a50aeae.html" target="_self">Mma&lt; gemm::GemmShape&lt; 8, 8, 16 &gt;, 32, uint8_t, layout::RowMajor, int8_t, layout::ColumnMajor, int, layout::RowMajor, OpMultiplyAdd &gt;</a></td><td class="desc">Matrix multiply-add operation: S32 = U8 * S8 + S32 </td></tr>
<tr id="row_0_0_27_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1arch_1_1Mma_3_01gemm_1_1GemmShape_3_018_00_018_00_0116_01_4_00_0132_00_01uint8_ab741d81fdc991345cb9e43c29fca573.html" target="_self">Mma&lt; gemm::GemmShape&lt; 8, 8, 16 &gt;, 32, uint8_t, layout::RowMajor, int8_t, layout::ColumnMajor, int, layout::RowMajor, OpMultiplyAddSaturate &gt;</a></td><td class="desc">Matrix multiply-add operation: S32 = U8 * S8 + S32 </td></tr>
<tr id="row_0_0_28_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1arch_1_1Mma_3_01gemm_1_1GemmShape_3_018_00_018_00_0116_01_4_00_0132_00_01uint8_5221708cec5828d35db1d1c47cb4964e.html" target="_self">Mma&lt; gemm::GemmShape&lt; 8, 8, 16 &gt;, 32, uint8_t, layout::RowMajor, uint8_t, layout::ColumnMajor, int, layout::RowMajor, OpMultiplyAdd &gt;</a></td><td class="desc">Matrix multiply-add operation: S32 = S8 * U8 + S32 </td></tr>
<tr id="row_0_0_29_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1arch_1_1Mma_3_01gemm_1_1GemmShape_3_018_00_018_00_0116_01_4_00_0132_00_01uint8_bef0c048bc0f8ba2d875cb7ab26d363b.html" target="_self">Mma&lt; gemm::GemmShape&lt; 8, 8, 16 &gt;, 32, uint8_t, layout::RowMajor, uint8_t, layout::ColumnMajor, int, layout::RowMajor, OpMultiplyAddSaturate &gt;</a></td><td class="desc">Matrix multiply-add operation: S32 = S8 * U8 + S32 </td></tr>
<tr id="row_0_0_30_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1arch_1_1Mma_3_01gemm_1_1GemmShape_3_018_00_018_00_0132_01_4_00_0132_00_01int4b_6e513ccbc44ae7909a60d93b9b5435b3.html" target="_self">Mma&lt; gemm::GemmShape&lt; 8, 8, 32 &gt;, 32, int4b_t, layout::RowMajor, int4b_t, layout::ColumnMajor, int, layout::RowMajor, OpMultiplyAdd &gt;</a></td><td class="desc">Matrix multiply-add operation: S32 = S4 * S4 + S32 </td></tr>
<tr id="row_0_0_31_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1arch_1_1Mma_3_01gemm_1_1GemmShape_3_018_00_018_00_0132_01_4_00_0132_00_01int4b_0ee08a4520882d24ba9026879265e892.html" target="_self">Mma&lt; gemm::GemmShape&lt; 8, 8, 32 &gt;, 32, int4b_t, layout::RowMajor, int4b_t, layout::ColumnMajor, int, layout::RowMajor, OpMultiplyAddSaturate &gt;</a></td><td class="desc">Matrix multiply-add operation: S32 = S4 * S4 + S32 </td></tr>
<tr id="row_0_0_32_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1arch_1_1Mma_3_01gemm_1_1GemmShape_3_018_00_018_00_0132_01_4_00_0132_00_01int4b_4746fc55e614df0016c518d3fda2677e.html" target="_self">Mma&lt; gemm::GemmShape&lt; 8, 8, 32 &gt;, 32, int4b_t, layout::RowMajor, uint4b_t, layout::ColumnMajor, int, layout::RowMajor, OpMultiplyAdd &gt;</a></td><td class="desc">Matrix multiply-add operation: S32 = S4 * U4 + S32 </td></tr>
<tr id="row_0_0_33_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1arch_1_1Mma_3_01gemm_1_1GemmShape_3_018_00_018_00_0132_01_4_00_0132_00_01int4b_546e9ec6de6a5970b326da6f6280f1d4.html" target="_self">Mma&lt; gemm::GemmShape&lt; 8, 8, 32 &gt;, 32, int4b_t, layout::RowMajor, uint4b_t, layout::ColumnMajor, int, layout::RowMajor, OpMultiplyAddSaturate &gt;</a></td><td class="desc">Matrix multiply-add operation: S32 = S4 * U4 + S32 </td></tr>
<tr id="row_0_0_34_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1arch_1_1Mma_3_01gemm_1_1GemmShape_3_018_00_018_00_0132_01_4_00_0132_00_01uint4b03e3b50dbcb30d0d1ac062f3a9d5abef.html" target="_self">Mma&lt; gemm::GemmShape&lt; 8, 8, 32 &gt;, 32, uint4b_t, layout::RowMajor, int4b_t, layout::ColumnMajor, int, layout::RowMajor, OpMultiplyAdd &gt;</a></td><td class="desc">Matrix multiply-add operation: S32 = U4 * S4 + S32 </td></tr>
<tr id="row_0_0_35_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1arch_1_1Mma_3_01gemm_1_1GemmShape_3_018_00_018_00_0132_01_4_00_0132_00_01uint4b6d968039dde5c9f062ab15f90a8049fe.html" target="_self">Mma&lt; gemm::GemmShape&lt; 8, 8, 32 &gt;, 32, uint4b_t, layout::RowMajor, int4b_t, layout::ColumnMajor, int, layout::RowMajor, OpMultiplyAddSaturate &gt;</a></td><td class="desc">Matrix multiply-add operation: S32 = U4 * S4 + S32 </td></tr>
<tr id="row_0_0_36_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1arch_1_1Mma_3_01gemm_1_1GemmShape_3_018_00_018_00_0132_01_4_00_0132_00_01uint4bc4b6ba004e25c44bfd9266c61f937dfb.html" target="_self">Mma&lt; gemm::GemmShape&lt; 8, 8, 32 &gt;, 32, uint4b_t, layout::RowMajor, uint4b_t, layout::ColumnMajor, int, layout::RowMajor, OpMultiplyAdd &gt;</a></td><td class="desc">Matrix multiply-add operation: S32 = U4 * U4 + S32 </td></tr>
<tr id="row_0_0_37_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1arch_1_1Mma_3_01gemm_1_1GemmShape_3_018_00_018_00_0132_01_4_00_0132_00_01uint4b451d5cf5d7e8cbbe476afe3dab5c09b2.html" target="_self">Mma&lt; gemm::GemmShape&lt; 8, 8, 32 &gt;, 32, uint4b_t, layout::RowMajor, uint4b_t, layout::ColumnMajor, int, layout::RowMajor, OpMultiplyAddSaturate &gt;</a></td><td class="desc">Matrix multiply-add operation: S32 = U4 * U4 + S32 </td></tr>
<tr id="row_0_0_38_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1arch_1_1Mma_3_01gemm_1_1GemmShape_3_018_00_018_00_014_01_4_00_018_00_01half__t_b0242d7a01097510effbc4718040d3e5.html" target="_self">Mma&lt; gemm::GemmShape&lt; 8, 8, 4 &gt;, 8, half_t, layout::ColumnMajor, half_t, layout::ColumnMajor, float, layout::RowMajor, OpMultiplyAdd &gt;</a></td><td class="desc">Matrix multiply-add operation: F32 = F16 * F16 + F32 </td></tr>
<tr id="row_0_0_39_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1arch_1_1Mma_3_01gemm_1_1GemmShape_3_018_00_018_00_014_01_4_00_018_00_01half__t_c7f88bfd32a544fba8111d2dcadeab11.html" target="_self">Mma&lt; gemm::GemmShape&lt; 8, 8, 4 &gt;, 8, half_t, layout::ColumnMajor, half_t, layout::ColumnMajor, half_t, layout::RowMajor, OpMultiplyAdd &gt;</a></td><td class="desc">Matrix multiply-add operation: F16 = F16 * F16 + F16 </td></tr>
<tr id="row_0_0_40_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1arch_1_1Mma_3_01gemm_1_1GemmShape_3_018_00_018_00_014_01_4_00_018_00_01half__t_44a3b2a8df88a2b067f1284515cb5371.html" target="_self">Mma&lt; gemm::GemmShape&lt; 8, 8, 4 &gt;, 8, half_t, layout::ColumnMajor, half_t, layout::RowMajor, float, layout::RowMajor, OpMultiplyAdd &gt;</a></td><td class="desc">Matrix multiply-add operation: F32 = F16 * F16 + F32 </td></tr>
<tr id="row_0_0_41_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1arch_1_1Mma_3_01gemm_1_1GemmShape_3_018_00_018_00_014_01_4_00_018_00_01half__t_4b7308177b308a272c1889fbe9670275.html" target="_self">Mma&lt; gemm::GemmShape&lt; 8, 8, 4 &gt;, 8, half_t, layout::ColumnMajor, half_t, layout::RowMajor, half_t, layout::RowMajor, OpMultiplyAdd &gt;</a></td><td class="desc">Matrix multiply-add operation: F16 = F16 * F16 + F16 </td></tr>
<tr id="row_0_0_42_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1arch_1_1Mma_3_01gemm_1_1GemmShape_3_018_00_018_00_014_01_4_00_018_00_01half__t_5a9888862cebd333ecaf11f7262f77d4.html" target="_self">Mma&lt; gemm::GemmShape&lt; 8, 8, 4 &gt;, 8, half_t, layout::RowMajor, half_t, layout::ColumnMajor, float, layout::RowMajor, OpMultiplyAdd &gt;</a></td><td class="desc">Matrix multiply-add operation: F32 = F16 * F16 + F32 </td></tr>
<tr id="row_0_0_43_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1arch_1_1Mma_3_01gemm_1_1GemmShape_3_018_00_018_00_014_01_4_00_018_00_01half__t_31defda8ea2b7d855642ffd77da1a411.html" target="_self">Mma&lt; gemm::GemmShape&lt; 8, 8, 4 &gt;, 8, half_t, layout::RowMajor, half_t, layout::ColumnMajor, half_t, layout::RowMajor, OpMultiplyAdd &gt;</a></td><td class="desc">Matrix multiply-add operation: F16 = F16 * F16 + F16 </td></tr>
<tr id="row_0_0_44_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1arch_1_1Mma_3_01gemm_1_1GemmShape_3_018_00_018_00_014_01_4_00_018_00_01half__t_839a7c8bb938d1661f4611e68f85d8cb.html" target="_self">Mma&lt; gemm::GemmShape&lt; 8, 8, 4 &gt;, 8, half_t, layout::RowMajor, half_t, layout::RowMajor, float, layout::RowMajor, OpMultiplyAdd &gt;</a></td><td class="desc">Matrix multiply-add operation: F32 = F16 * F16 + F32 </td></tr>
<tr id="row_0_0_45_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1arch_1_1Mma_3_01gemm_1_1GemmShape_3_018_00_018_00_014_01_4_00_018_00_01half__t_73d9802d6b944a5299bc255887db6bbc.html" target="_self">Mma&lt; gemm::GemmShape&lt; 8, 8, 4 &gt;, 8, half_t, layout::RowMajor, half_t, layout::RowMajor, half_t, layout::RowMajor, OpMultiplyAdd &gt;</a></td><td class="desc">Matrix multiply-add operation: F16 = F16 * F16 + F16 </td></tr>
<tr id="row_0_0_46_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1arch_1_1PtxWmma.html" target="_self">PtxWmma</a></td><td class="desc">WMMA Matrix multiply-add operation </td></tr>
<tr id="row_0_0_47_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1arch_1_1PtxWmmaLoadA.html" target="_self">PtxWmmaLoadA</a></td><td class="desc">WMMA PTX string load for A, B, and C matrices </td></tr>
<tr id="row_0_0_48_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1arch_1_1PtxWmmaLoadB.html" target="_self">PtxWmmaLoadB</a></td><td class="desc"></td></tr>
<tr id="row_0_0_49_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1arch_1_1PtxWmmaLoadC.html" target="_self">PtxWmmaLoadC</a></td><td class="desc"></td></tr>
<tr id="row_0_0_50_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1arch_1_1PtxWmmaStoreD.html" target="_self">PtxWmmaStoreD</a></td><td class="desc">WMMA store for matrix D </td></tr>
<tr id="row_0_0_51_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1arch_1_1Sm50.html" target="_self">Sm50</a></td><td class="desc"></td></tr>
<tr id="row_0_0_52_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1arch_1_1Sm60.html" target="_self">Sm60</a></td><td class="desc"></td></tr>
<tr id="row_0_0_53_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1arch_1_1Sm61.html" target="_self">Sm61</a></td><td class="desc"></td></tr>
<tr id="row_0_0_54_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1arch_1_1Sm70.html" target="_self">Sm70</a></td><td class="desc"></td></tr>
<tr id="row_0_0_55_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1arch_1_1Sm72.html" target="_self">Sm72</a></td><td class="desc"></td></tr>
<tr id="row_0_0_56_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1arch_1_1Sm75.html" target="_self">Sm75</a></td><td class="desc"></td></tr>
<tr id="row_0_0_57_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1arch_1_1Wmma_3_01Shape___00_01cutlass_1_1half__t_00_01LayoutA___00_01cutlass_1_84e30c8cc93eeb7ca02f651bd16d4c38.html" target="_self">Wmma&lt; Shape_, cutlass::half_t, LayoutA_, cutlass::half_t, LayoutB_, ElementC_, LayoutC_, cutlass::arch::OpMultiplyAdd &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_0_58_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1arch_1_1Wmma_3_01Shape___00_01cutlass_1_1int4b__t_00_01LayoutA___00_01cutlass_16fd808a90b3cf9d7cfc99f30888ca3fe.html" target="_self">Wmma&lt; Shape_, cutlass::int4b_t, LayoutA_, cutlass::int4b_t, LayoutB_, int32_t, LayoutC_, cutlass::arch::OpMultiplyAdd &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_0_59_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1arch_1_1Wmma_3_01Shape___00_01cutlass_1_1uint1b__t_00_01LayoutA___00_01cutlass_c80a7ea4d219cd9b13b560b493338028.html" target="_self">Wmma&lt; Shape_, cutlass::uint1b_t, LayoutA_, cutlass::uint1b_t, LayoutB_, int32_t, LayoutC_, cutlass::arch::OpXorPopc &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_0_60_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1arch_1_1Wmma_3_01Shape___00_01int8__t_00_01LayoutA___00_01int8__t_00_01LayoutB_505c57bb6818a941dc16f00cf35a9ec0.html" target="_self">Wmma&lt; Shape_, int8_t, LayoutA_, int8_t, LayoutB_, int32_t, LayoutC_, cutlass::arch::OpMultiplyAdd &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_0_61_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1arch_1_1Wmma_3_01Shape___00_01uint8__t_00_01LayoutA___00_01uint8__t_00_01Layout219a464a1248ebfc37aa29bcb10cb1b0.html" target="_self">Wmma&lt; Shape_, uint8_t, LayoutA_, uint8_t, LayoutB_, int32_t, LayoutC_, cutlass::arch::OpMultiplyAdd &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_1_" style="display:none;"><td class="entry"><span style="width:16px;display:inline-block;">&#160;</span><span id="arr_0_1_" class="arrow" onclick="toggleFolder('0_1_')">&#9658;</span><span class="icona"><span class="icon">N</span></span><a class="el" href="namespacecutlass_1_1device__memory.html" target="_self">device_memory</a></td><td class="desc"></td></tr>
<tr id="row_0_1_0_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span id="arr_0_1_0_" class="arrow" onclick="toggleFolder('0_1_0_')">&#9658;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1device__memory_1_1allocation.html" target="_self">allocation</a></td><td class="desc">Device allocation abstraction that tracks size and capacity </td></tr>
<tr id="row_0_1_0_0_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1device__memory_1_1allocation_1_1deleter.html" target="_self">deleter</a></td><td class="desc">Delete functor for CUDA device memory </td></tr>
<tr id="row_0_2_" style="display:none;"><td class="entry"><span style="width:16px;display:inline-block;">&#160;</span><span id="arr_0_2_" class="arrow" onclick="toggleFolder('0_2_')">&#9658;</span><span class="icona"><span class="icon">N</span></span><a class="el" href="namespacecutlass_1_1epilogue.html" target="_self">epilogue</a></td><td class="desc"></td></tr>
<tr id="row_0_2_0_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span id="arr_0_2_0_" class="arrow" onclick="toggleFolder('0_2_0_')">&#9658;</span><span class="icona"><span class="icon">N</span></span><a class="el" href="namespacecutlass_1_1epilogue_1_1thread.html" target="_self">thread</a></td><td class="desc"></td></tr>
<tr id="row_0_2_0_0_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span id="arr_0_2_0_0_" class="arrow" onclick="toggleFolder('0_2_0_0_')">&#9658;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1epilogue_1_1thread_1_1Convert.html" target="_self">Convert</a></td><td class="desc"></td></tr>
<tr id="row_0_2_0_0_0_" style="display:none;"><td class="entry"><span style="width:80px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1epilogue_1_1thread_1_1Convert_1_1Params.html" target="_self">Params</a></td><td class="desc">Host-constructable parameters structure </td></tr>
<tr id="row_0_2_0_1_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span id="arr_0_2_0_1_" class="arrow" onclick="toggleFolder('0_2_0_1_')">&#9658;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1epilogue_1_1thread_1_1LinearCombination.html" target="_self">LinearCombination</a></td><td class="desc"></td></tr>
<tr id="row_0_2_0_1_0_" style="display:none;"><td class="entry"><span style="width:80px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1epilogue_1_1thread_1_1LinearCombination_1_1Params.html" target="_self">Params</a></td><td class="desc">Host-constructable parameters structure </td></tr>
<tr id="row_0_2_0_2_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span id="arr_0_2_0_2_" class="arrow" onclick="toggleFolder('0_2_0_2_')">&#9658;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1epilogue_1_1thread_1_1LinearCombinationClamp.html" target="_self">LinearCombinationClamp</a></td><td class="desc"></td></tr>
<tr id="row_0_2_0_2_0_" style="display:none;"><td class="entry"><span style="width:80px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1epilogue_1_1thread_1_1LinearCombinationClamp_1_1Params.html" target="_self">Params</a></td><td class="desc">Host-constructable parameters structure </td></tr>
<tr id="row_0_2_0_3_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span id="arr_0_2_0_3_" class="arrow" onclick="toggleFolder('0_2_0_3_')">&#9658;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1epilogue_1_1thread_1_1LinearCombinationRelu.html" target="_self">LinearCombinationRelu</a></td><td class="desc"></td></tr>
<tr id="row_0_2_0_3_0_" style="display:none;"><td class="entry"><span style="width:80px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1epilogue_1_1thread_1_1LinearCombinationRelu_1_1Params.html" target="_self">Params</a></td><td class="desc">Host-constructable parameters structure </td></tr>
<tr id="row_0_2_0_4_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span id="arr_0_2_0_4_" class="arrow" onclick="toggleFolder('0_2_0_4_')">&#9658;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1epilogue_1_1thread_1_1LinearCombinationRelu_3_01ElementOutput___00_01Count_00_01int_00_01float_00_01Round_01_4.html" target="_self">LinearCombinationRelu&lt; ElementOutput_, Count, int, float, Round &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_2_0_4_0_" style="display:none;"><td class="entry"><span style="width:80px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1epilogue_1_1thread_1_1LinearCombinationRelu_3_01ElementOutput___00_01Count_00_00274a94522c46cd041d0b10d484e2ef3.html" target="_self">Params</a></td><td class="desc">Host-constructable parameters structure </td></tr>
<tr id="row_0_2_0_5_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span id="arr_0_2_0_5_" class="arrow" onclick="toggleFolder('0_2_0_5_')">&#9658;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1epilogue_1_1thread_1_1ReductionOpPlus.html" target="_self">ReductionOpPlus</a></td><td class="desc"></td></tr>
<tr id="row_0_2_0_5_0_" style="display:none;"><td class="entry"><span style="width:80px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1epilogue_1_1thread_1_1ReductionOpPlus_1_1Params.html" target="_self">Params</a></td><td class="desc">Host-constructable parameters structure </td></tr>
<tr id="row_0_2_1_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span id="arr_0_2_1_" class="arrow" onclick="toggleFolder('0_2_1_')">&#9658;</span><span class="icona"><span class="icon">N</span></span><a class="el" href="namespacecutlass_1_1epilogue_1_1threadblock.html" target="_self">threadblock</a></td><td class="desc"></td></tr>
<tr id="row_0_2_1_0_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span id="arr_0_2_1_0_" class="arrow" onclick="toggleFolder('0_2_1_0_')">&#9658;</span><span class="icona"><span class="icon">N</span></span><a class="el" href="namespacecutlass_1_1epilogue_1_1threadblock_1_1detail.html" target="_self">detail</a></td><td class="desc"></td></tr>
<tr id="row_0_2_1_0_0_" style="display:none;"><td class="entry"><span style="width:80px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1epilogue_1_1threadblock_1_1detail_1_1RowArrangement.html" target="_self">RowArrangement</a></td><td class="desc"><a class="el" href="structcutlass_1_1epilogue_1_1threadblock_1_1detail_1_1RowArrangement.html" title="RowArrangement determines how one or more warps cover a region of consecutive rows. ">RowArrangement</a> determines how one or more warps cover a region of consecutive rows </td></tr>
<tr id="row_0_2_1_0_1_" style="display:none;"><td class="entry"><span style="width:80px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1epilogue_1_1threadblock_1_1detail_1_1RowArrangement_3_01Shape_00_01WarpsRemaini91159e6f7e123d881e3ec45101fa4f81.html" target="_self">RowArrangement&lt; Shape, WarpsRemaining, ElementsPerAccess, ElementSize, false &gt;</a></td><td class="desc"><a class="el" href="structcutlass_1_1epilogue_1_1threadblock_1_1detail_1_1RowArrangement.html" title="RowArrangement determines how one or more warps cover a region of consecutive rows. ">RowArrangement</a> in which each warp's access is a 1D tiled arrangement </td></tr>
<tr id="row_0_2_1_0_2_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span id="arr_0_2_1_0_2_" class="arrow" onclick="toggleFolder('0_2_1_0_2_')">&#9658;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1epilogue_1_1threadblock_1_1detail_1_1RowArrangement_3_01Shape_00_01WarpsRemaini6d8790249bf12cac580da73bb37eb791.html" target="_self">RowArrangement&lt; Shape, WarpsRemaining, ElementsPerAccess, ElementSize, true &gt;</a></td><td class="desc"><a class="el" href="structcutlass_1_1epilogue_1_1threadblock_1_1detail_1_1RowArrangement.html" title="RowArrangement determines how one or more warps cover a region of consecutive rows. ">RowArrangement</a> in which each warp's access is a 2D tiled arrangement </td></tr>
<tr id="row_0_2_1_0_2_0_" style="display:none;"><td class="entry"><span style="width:96px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1epilogue_1_1threadblock_1_1detail_1_1RowArrangement_3_01Shape_00_01WarpsRemainief28e98b3f284469f271d28aba73de2e.html" target="_self">Detail</a></td><td class="desc"></td></tr>
<tr id="row_0_2_1_1_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1epilogue_1_1threadblock_1_1DefaultEpilogueComplexTensorOp.html" target="_self">DefaultEpilogueComplexTensorOp</a></td><td class="desc">Defines sensible defaults for epilogues for TensorOps </td></tr>
<tr id="row_0_2_1_2_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1epilogue_1_1threadblock_1_1DefaultEpilogueSimt.html" target="_self">DefaultEpilogueSimt</a></td><td class="desc">Defines sensible defaults for epilogues for SimtOps </td></tr>
<tr id="row_0_2_1_3_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1epilogue_1_1threadblock_1_1DefaultEpilogueTensorOp.html" target="_self">DefaultEpilogueTensorOp</a></td><td class="desc">Defines sensible defaults for epilogues for TensorOps </td></tr>
<tr id="row_0_2_1_4_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1epilogue_1_1threadblock_1_1DefaultEpilogueVoltaTensorOp.html" target="_self">DefaultEpilogueVoltaTensorOp</a></td><td class="desc">Defines sensible defaults for epilogues for TensorOps </td></tr>
<tr id="row_0_2_1_5_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1epilogue_1_1threadblock_1_1DefaultEpilogueWmmaTensorOp.html" target="_self">DefaultEpilogueWmmaTensorOp</a></td><td class="desc">Defines sensible defaults for epilogues for WMMA TensorOps </td></tr>
<tr id="row_0_2_1_6_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1epilogue_1_1threadblock_1_1DefaultInterleavedEpilogueTensorOp.html" target="_self">DefaultInterleavedEpilogueTensorOp</a></td><td class="desc"></td></tr>
<tr id="row_0_2_1_7_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span id="arr_0_2_1_7_" class="arrow" onclick="toggleFolder('0_2_1_7_')">&#9658;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1epilogue_1_1threadblock_1_1DefaultInterleavedThreadMapTensorOp.html" target="_self">DefaultInterleavedThreadMapTensorOp</a></td><td class="desc">Defines the optimal thread map for TensorOp accumulator layouts </td></tr>
<tr id="row_0_2_1_7_0_" style="display:none;"><td class="entry"><span style="width:80px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1epilogue_1_1threadblock_1_1DefaultInterleavedThreadMapTensorOp_1_1Detail.html" target="_self">Detail</a></td><td class="desc"></td></tr>
<tr id="row_0_2_1_8_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span id="arr_0_2_1_8_" class="arrow" onclick="toggleFolder('0_2_1_8_')">&#9658;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1epilogue_1_1threadblock_1_1DefaultThreadMapSimt.html" target="_self">DefaultThreadMapSimt</a></td><td class="desc">Defines the optimal thread map for SIMT accumulator layouts </td></tr>
<tr id="row_0_2_1_8_0_" style="display:none;"><td class="entry"><span style="width:80px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1epilogue_1_1threadblock_1_1DefaultThreadMapSimt_1_1Detail.html" target="_self">Detail</a></td><td class="desc"></td></tr>
<tr id="row_0_2_1_9_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span id="arr_0_2_1_9_" class="arrow" onclick="toggleFolder('0_2_1_9_')">&#9658;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1epilogue_1_1threadblock_1_1DefaultThreadMapTensorOp.html" target="_self">DefaultThreadMapTensorOp</a></td><td class="desc">Defines the optimal thread map for TensorOp accumulator layouts </td></tr>
<tr id="row_0_2_1_9_0_" style="display:none;"><td class="entry"><span style="width:80px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1epilogue_1_1threadblock_1_1DefaultThreadMapTensorOp_1_1Detail.html" target="_self">Detail</a></td><td class="desc"></td></tr>
<tr id="row_0_2_1_10_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1epilogue_1_1threadblock_1_1DefaultThreadMapVoltaTensorOp.html" target="_self">DefaultThreadMapVoltaTensorOp</a></td><td class="desc">Defines the optimal thread map for TensorOp accumulator layouts </td></tr>
<tr id="row_0_2_1_11_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span id="arr_0_2_1_11_" class="arrow" onclick="toggleFolder('0_2_1_11_')">&#9658;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1epilogue_1_1threadblock_1_1DefaultThreadMapVoltaTensorOp_3_01ThreadblockShape__95db04b7b72e34283958bd7fbf851d16.html" target="_self">DefaultThreadMapVoltaTensorOp&lt; ThreadblockShape_, WarpShape_, PartitionsK, ElementOutput_, ElementsPerAccess, float &gt;</a></td><td class="desc">Defines the optimal thread map for TensorOp accumulator layouts </td></tr>
<tr id="row_0_2_1_11_0_" style="display:none;"><td class="entry"><span style="width:80px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1epilogue_1_1threadblock_1_1DefaultThreadMapVoltaTensorOp_3_01ThreadblockShape__52116c60c62f0fd520071558e42b814f.html" target="_self">Detail</a></td><td class="desc"></td></tr>
<tr id="row_0_2_1_12_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span id="arr_0_2_1_12_" class="arrow" onclick="toggleFolder('0_2_1_12_')">&#9658;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1epilogue_1_1threadblock_1_1DefaultThreadMapVoltaTensorOp_3_01ThreadblockShape__d58c94abc36b7c5c109b55202c6992e7.html" target="_self">DefaultThreadMapVoltaTensorOp&lt; ThreadblockShape_, WarpShape_, PartitionsK, ElementOutput_, ElementsPerAccess, half_t &gt;</a></td><td class="desc">Defines the optimal thread map for TensorOp accumulator layouts </td></tr>
<tr id="row_0_2_1_12_0_" style="display:none;"><td class="entry"><span style="width:80px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1epilogue_1_1threadblock_1_1DefaultThreadMapVoltaTensorOp_3_01ThreadblockShape__4433cc988100e98097a748d2670fb0fc.html" target="_self">Detail</a></td><td class="desc"></td></tr>
<tr id="row_0_2_1_13_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span id="arr_0_2_1_13_" class="arrow" onclick="toggleFolder('0_2_1_13_')">&#9658;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1epilogue_1_1threadblock_1_1DefaultThreadMapWmmaTensorOp.html" target="_self">DefaultThreadMapWmmaTensorOp</a></td><td class="desc">Defines the optimal thread map for Wmma TensorOp accumulator layouts </td></tr>
<tr id="row_0_2_1_13_0_" style="display:none;"><td class="entry"><span style="width:80px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1epilogue_1_1threadblock_1_1DefaultThreadMapWmmaTensorOp_1_1Detail.html" target="_self">Detail</a></td><td class="desc"></td></tr>
<tr id="row_0_2_1_14_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span id="arr_0_2_1_14_" class="arrow" onclick="toggleFolder('0_2_1_14_')">&#9658;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1epilogue_1_1threadblock_1_1DirectEpilogueTensorOp.html" target="_self">DirectEpilogueTensorOp</a></td><td class="desc"><a class="el" href="classcutlass_1_1epilogue_1_1threadblock_1_1Epilogue.html" title="Epilogue operator without splitk. ">Epilogue</a> operator </td></tr>
<tr id="row_0_2_1_14_0_" style="display:none;"><td class="entry"><span style="width:80px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1epilogue_1_1threadblock_1_1DirectEpilogueTensorOp_1_1Params.html" target="_self">Params</a></td><td class="desc">Parameters structure for host-constructible state </td></tr>
<tr id="row_0_2_1_14_1_" style="display:none;"><td class="entry"><span style="width:80px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1epilogue_1_1threadblock_1_1DirectEpilogueTensorOp_1_1SharedStorage.html" target="_self">SharedStorage</a></td><td class="desc">Shared storage allocation needed by the epilogue </td></tr>
<tr id="row_0_2_1_15_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1epilogue_1_1threadblock_1_1Epilogue.html" target="_self">Epilogue</a></td><td class="desc"><a class="el" href="classcutlass_1_1epilogue_1_1threadblock_1_1Epilogue.html" title="Epilogue operator without splitk. ">Epilogue</a> operator without splitk </td></tr>
<tr id="row_0_2_1_16_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span id="arr_0_2_1_16_" class="arrow" onclick="toggleFolder('0_2_1_16_')">&#9658;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1epilogue_1_1threadblock_1_1EpilogueBase.html" target="_self">EpilogueBase</a></td><td class="desc">Base class for epilogues defining warp-level </td></tr>
<tr id="row_0_2_1_16_0_" style="display:none;"><td class="entry"><span style="width:80px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1epilogue_1_1threadblock_1_1EpilogueBase_1_1SharedStorage.html" target="_self">SharedStorage</a></td><td class="desc">Shared storage allocation needed by the epilogue </td></tr>
<tr id="row_0_2_1_17_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span id="arr_0_2_1_17_" class="arrow" onclick="toggleFolder('0_2_1_17_')">&#9658;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1epilogue_1_1threadblock_1_1InterleavedEpilogue.html" target="_self">InterleavedEpilogue</a></td><td class="desc"><a class="el" href="classcutlass_1_1epilogue_1_1threadblock_1_1Epilogue.html" title="Epilogue operator without splitk. ">Epilogue</a> operator without splitk </td></tr>
<tr id="row_0_2_1_17_0_" style="display:none;"><td class="entry"><span style="width:80px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1epilogue_1_1threadblock_1_1InterleavedEpilogue_1_1SharedStorage.html" target="_self">SharedStorage</a></td><td class="desc">Shared storage allocation needed by the epilogue </td></tr>
<tr id="row_0_2_1_18_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span id="arr_0_2_1_18_" class="arrow" onclick="toggleFolder('0_2_1_18_')">&#9658;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1epilogue_1_1threadblock_1_1InterleavedOutputTileThreadMap.html" target="_self">InterleavedOutputTileThreadMap</a></td><td class="desc"></td></tr>
<tr id="row_0_2_1_18_0_" style="display:none;"><td class="entry"><span style="width:80px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1epilogue_1_1threadblock_1_1InterleavedOutputTileThreadMap_1_1Detail.html" target="_self">Detail</a></td><td class="desc"></td></tr>
<tr id="row_0_2_1_19_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span id="arr_0_2_1_19_" class="arrow" onclick="toggleFolder('0_2_1_19_')">&#9658;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1epilogue_1_1threadblock_1_1InterleavedPredicatedTileIterator.html" target="_self">InterleavedPredicatedTileIterator</a></td><td class="desc"></td></tr>
<tr id="row_0_2_1_19_0_" style="display:none;"><td class="entry"><span style="width:80px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1epilogue_1_1threadblock_1_1InterleavedPredicatedTileIterator_1_1Mask.html" target="_self">Mask</a></td><td class="desc"><a class="el" href="structcutlass_1_1epilogue_1_1threadblock_1_1InterleavedPredicatedTileIterator_1_1Mask.html" title="Mask object. ">Mask</a> object </td></tr>
<tr id="row_0_2_1_19_1_" style="display:none;"><td class="entry"><span style="width:80px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1epilogue_1_1threadblock_1_1InterleavedPredicatedTileIterator_1_1Params.html" target="_self">Params</a></td><td class="desc"></td></tr>
<tr id="row_0_2_1_20_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span id="arr_0_2_1_20_" class="arrow" onclick="toggleFolder('0_2_1_20_')">&#9658;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1epilogue_1_1threadblock_1_1OutputTileOptimalThreadMap.html" target="_self">OutputTileOptimalThreadMap</a></td><td class="desc"></td></tr>
<tr id="row_0_2_1_20_0_" style="display:none;"><td class="entry"><span style="width:80px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1epilogue_1_1threadblock_1_1OutputTileOptimalThreadMap_1_1CompactedThreadMap.html" target="_self">CompactedThreadMap</a></td><td class="desc">Compacted thread map in which the 4D region is contiguous </td></tr>
<tr id="row_0_2_1_20_1_" style="display:none;"><td class="entry"><span style="width:80px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1epilogue_1_1threadblock_1_1OutputTileOptimalThreadMap_1_1Detail.html" target="_self">Detail</a></td><td class="desc"></td></tr>
<tr id="row_0_2_1_21_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1epilogue_1_1threadblock_1_1OutputTileShape.html" target="_self">OutputTileShape</a></td><td class="desc">Tuple defining point in output tile </td></tr>
<tr id="row_0_2_1_22_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1epilogue_1_1threadblock_1_1OutputTileThreadMap.html" target="_self">OutputTileThreadMap</a></td><td class="desc"></td></tr>
<tr id="row_0_2_1_23_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span id="arr_0_2_1_23_" class="arrow" onclick="toggleFolder('0_2_1_23_')">&#9658;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1epilogue_1_1threadblock_1_1PredicatedTileIterator.html" target="_self">PredicatedTileIterator</a></td><td class="desc"></td></tr>
<tr id="row_0_2_1_23_0_" style="display:none;"><td class="entry"><span style="width:80px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1epilogue_1_1threadblock_1_1PredicatedTileIterator_1_1Mask.html" target="_self">Mask</a></td><td class="desc"><a class="el" href="structcutlass_1_1epilogue_1_1threadblock_1_1PredicatedTileIterator_1_1Mask.html" title="Mask object. ">Mask</a> object </td></tr>
<tr id="row_0_2_1_23_1_" style="display:none;"><td class="entry"><span style="width:80px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1epilogue_1_1threadblock_1_1PredicatedTileIterator_1_1Params.html" target="_self">Params</a></td><td class="desc"></td></tr>
<tr id="row_0_2_1_24_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1epilogue_1_1threadblock_1_1SharedLoadIterator.html" target="_self">SharedLoadIterator</a></td><td class="desc"></td></tr>
<tr id="row_0_2_2_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span id="arr_0_2_2_" class="arrow" onclick="toggleFolder('0_2_2_')">&#9658;</span><span class="icona"><span class="icon">N</span></span><a class="el" href="namespacecutlass_1_1epilogue_1_1warp.html" target="_self">warp</a></td><td class="desc"></td></tr>
<tr id="row_0_2_2_0_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1epilogue_1_1warp_1_1FragmentIteratorComplexTensorOp.html" target="_self">FragmentIteratorComplexTensorOp</a></td><td class="desc"></td></tr>
<tr id="row_0_2_2_1_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1epilogue_1_1warp_1_1FragmentIteratorComplexTensorOp_3_01WarpShape___00_01Operato8cf03c624cf3210c71b7cbd580b080f8.html" target="_self">FragmentIteratorComplexTensorOp&lt; WarpShape_, OperatorShape_, OperatorElementC_, OperatorFragmentC_, layout::RowMajor &gt;</a></td><td class="desc">Partial specialization for row-major shared memory </td></tr>
<tr id="row_0_2_2_2_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1epilogue_1_1warp_1_1FragmentIteratorSimt.html" target="_self">FragmentIteratorSimt</a></td><td class="desc">Fragment iterator for SIMT accumulator arrangements </td></tr>
<tr id="row_0_2_2_3_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1epilogue_1_1warp_1_1FragmentIteratorSimt_3_01WarpShape___00_01Operator___00_01la3f2abc523201c1b0228df99119ab88e1.html" target="_self">FragmentIteratorSimt&lt; WarpShape_, Operator_, layout::RowMajor, MmaSimtPolicy_ &gt;</a></td><td class="desc">Partial specialization for row-major shared memory </td></tr>
<tr id="row_0_2_2_4_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1epilogue_1_1warp_1_1FragmentIteratorTensorOp.html" target="_self">FragmentIteratorTensorOp</a></td><td class="desc"></td></tr>
<tr id="row_0_2_2_5_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1epilogue_1_1warp_1_1FragmentIteratorTensorOp_3_01WarpShape___00_01OperatorShape_e459aab140a2ce78336e584f95886726.html" target="_self">FragmentIteratorTensorOp&lt; WarpShape_, OperatorShape_, OperatorElementC_, OperatorFragmentC_, layout::ColumnMajorInterleaved&lt; InterleavedK &gt; &gt;</a></td><td class="desc">Dedicated to interleaved layout </td></tr>
<tr id="row_0_2_2_6_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1epilogue_1_1warp_1_1FragmentIteratorTensorOp_3_01WarpShape___00_01OperatorShape_5e78dabe303f20d76b00c600aab61eda.html" target="_self">FragmentIteratorTensorOp&lt; WarpShape_, OperatorShape_, OperatorElementC_, OperatorFragmentC_, layout::RowMajor &gt;</a></td><td class="desc">Partial specialization for row-major shared memory </td></tr>
<tr id="row_0_2_2_7_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1epilogue_1_1warp_1_1FragmentIteratorVoltaTensorOp.html" target="_self">FragmentIteratorVoltaTensorOp</a></td><td class="desc"></td></tr>
<tr id="row_0_2_2_8_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1epilogue_1_1warp_1_1FragmentIteratorVoltaTensorOp_3_01WarpShape___00_01gemm_1_1Gdb805a2dc5571ac3b66e0fe6ffdcede2.html" target="_self">FragmentIteratorVoltaTensorOp&lt; WarpShape_, gemm::GemmShape&lt; 32, 32, 4 &gt;, float, layout::RowMajor &gt;</a></td><td class="desc">Partial specialization for row-major shared memory </td></tr>
<tr id="row_0_2_2_9_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1epilogue_1_1warp_1_1FragmentIteratorVoltaTensorOp_3_01WarpShape___00_01gemm_1_1G16e08718cffa0989cce3fe8dbc4b075b.html" target="_self">FragmentIteratorVoltaTensorOp&lt; WarpShape_, gemm::GemmShape&lt; 32, 32, 4 &gt;, half_t, layout::RowMajor &gt;</a></td><td class="desc">Partial specialization for row-major shared memory </td></tr>
<tr id="row_0_2_2_10_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1epilogue_1_1warp_1_1FragmentIteratorWmmaTensorOp.html" target="_self">FragmentIteratorWmmaTensorOp</a></td><td class="desc"></td></tr>
<tr id="row_0_2_2_11_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1epilogue_1_1warp_1_1FragmentIteratorWmmaTensorOp_3_01WarpShape___00_01OperatorShfdb1f120c6797383663f9fd11d0fc599.html" target="_self">FragmentIteratorWmmaTensorOp&lt; WarpShape_, OperatorShape_, OperatorElementC_, OperatorFragmentC_, layout::RowMajor &gt;</a></td><td class="desc">Partial specialization for row-major shared memory </td></tr>
<tr id="row_0_2_2_12_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1epilogue_1_1warp_1_1SimtPolicy.html" target="_self">SimtPolicy</a></td><td class="desc"></td></tr>
<tr id="row_0_2_2_13_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1epilogue_1_1warp_1_1SimtPolicy_3_01WarpShape___00_01Operator___00_01layout_1_1Rcef1c60e23e997017ae176c92931151d.html" target="_self">SimtPolicy&lt; WarpShape_, Operator_, layout::RowMajor, MmaSimtPolicy_ &gt;</a></td><td class="desc">Partial specialization for row-major </td></tr>
<tr id="row_0_2_2_14_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1epilogue_1_1warp_1_1TensorOpPolicy.html" target="_self">TensorOpPolicy</a></td><td class="desc">Policy details related to the epilogue </td></tr>
<tr id="row_0_2_2_15_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1epilogue_1_1warp_1_1TensorOpPolicy_3_01WarpShape_00_01OperatorShape_00_01layout69549d10c3610d943987eb90e827bc05.html" target="_self">TensorOpPolicy&lt; WarpShape, OperatorShape, layout::ColumnMajorInterleaved&lt; InterleavedK &gt; &gt;</a></td><td class="desc">Partial specialization for column-major-interleaved </td></tr>
<tr id="row_0_2_2_16_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1epilogue_1_1warp_1_1TensorOpPolicy_3_01WarpShape_00_01OperatorShape_00_01layout_1_1RowMajor_01_4.html" target="_self">TensorOpPolicy&lt; WarpShape, OperatorShape, layout::RowMajor &gt;</a></td><td class="desc">Partial specialization for row-major </td></tr>
<tr id="row_0_2_2_17_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1epilogue_1_1warp_1_1TileIteratorSimt.html" target="_self">TileIteratorSimt</a></td><td class="desc">Template for reading and writing tiles of accumulators to shared memory </td></tr>
<tr id="row_0_2_2_18_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1epilogue_1_1warp_1_1TileIteratorSimt_3_01WarpShape___00_01Operator___00_01Elemenf2bd262ed3e202b25d5802d83965bf3b.html" target="_self">TileIteratorSimt&lt; WarpShape_, Operator_, Element_, layout::RowMajor, MmaSimtPolicy_ &gt;</a></td><td class="desc">Template for reading and writing tiles of accumulators to shared memory </td></tr>
<tr id="row_0_2_2_19_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1epilogue_1_1warp_1_1TileIteratorTensorOp.html" target="_self">TileIteratorTensorOp</a></td><td class="desc">Template for reading and writing tiles of accumulators to shared memory </td></tr>
<tr id="row_0_2_2_20_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span id="arr_0_2_2_20_" class="arrow" onclick="toggleFolder('0_2_2_20_')">&#9658;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1epilogue_1_1warp_1_1TileIteratorTensorOp_3_01WarpShape___00_01OperatorShape___003cbb32beb84b4984cb7853662096d289.html" target="_self">TileIteratorTensorOp&lt; WarpShape_, OperatorShape_, Element_, layout::RowMajor &gt;</a></td><td class="desc">Template for reading and writing tiles of accumulators to shared memory </td></tr>
<tr id="row_0_2_2_20_0_" style="display:none;"><td class="entry"><span style="width:80px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1epilogue_1_1warp_1_1TileIteratorTensorOp_3_01WarpShape___00_01OperatorShape___05f11e023c9e6ee5f7a888fa4c5bbf6d1.html" target="_self">Detail</a></td><td class="desc"></td></tr>
<tr id="row_0_2_2_21_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1epilogue_1_1warp_1_1TileIteratorVoltaTensorOp.html" target="_self">TileIteratorVoltaTensorOp</a></td><td class="desc">Template for reading and writing tiles of accumulators to shared memory </td></tr>
<tr id="row_0_2_2_22_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span id="arr_0_2_2_22_" class="arrow" onclick="toggleFolder('0_2_2_22_')">&#9658;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1epilogue_1_1warp_1_1TileIteratorVoltaTensorOp_3_01WarpShape___00_01gemm_1_1GemmS2fe0c60b727c738c622c18fc3dd76644.html" target="_self">TileIteratorVoltaTensorOp&lt; WarpShape_, gemm::GemmShape&lt; 32, 32, 4 &gt;, float, layout::RowMajor &gt;</a></td><td class="desc">Template for reading and writing tiles of accumulators to shared memory </td></tr>
<tr id="row_0_2_2_22_0_" style="display:none;"><td class="entry"><span style="width:80px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1epilogue_1_1warp_1_1TileIteratorVoltaTensorOp_3_01WarpShape___00_01gemm_1_1Gemm770cbca45441d295d5d7433e8222a700.html" target="_self">Detail</a></td><td class="desc"></td></tr>
<tr id="row_0_2_2_23_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span id="arr_0_2_2_23_" class="arrow" onclick="toggleFolder('0_2_2_23_')">&#9658;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1epilogue_1_1warp_1_1TileIteratorVoltaTensorOp_3_01WarpShape___00_01gemm_1_1GemmSa0ceeeddc22575876eb977da7f5416a8.html" target="_self">TileIteratorVoltaTensorOp&lt; WarpShape_, gemm::GemmShape&lt; 32, 32, 4 &gt;, half_t, layout::RowMajor &gt;</a></td><td class="desc">Template for reading and writing tiles of accumulators to shared memory </td></tr>
<tr id="row_0_2_2_23_0_" style="display:none;"><td class="entry"><span style="width:80px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1epilogue_1_1warp_1_1TileIteratorVoltaTensorOp_3_01WarpShape___00_01gemm_1_1Gemmffcab2297c8de8d0013602a39c525b78.html" target="_self">Detail</a></td><td class="desc"></td></tr>
<tr id="row_0_2_2_24_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1epilogue_1_1warp_1_1TileIteratorWmmaTensorOp.html" target="_self">TileIteratorWmmaTensorOp</a></td><td class="desc">Template for reading and writing tiles of accumulators to shared memory </td></tr>
<tr id="row_0_2_2_25_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1epilogue_1_1warp_1_1TileIteratorWmmaTensorOp_3_01WarpShape___00_01OperatorShape_fd6a91cd8bbd07ecd1344326b830e3a4.html" target="_self">TileIteratorWmmaTensorOp&lt; WarpShape_, OperatorShape_, OperatorFragment_, layout::RowMajor &gt;</a></td><td class="desc">Template for reading and writing tiles of accumulators to shared memory </td></tr>
<tr id="row_0_2_2_26_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1epilogue_1_1warp_1_1VoltaTensorOpPolicy.html" target="_self">VoltaTensorOpPolicy</a></td><td class="desc">Policy details related to the epilogue </td></tr>
<tr id="row_0_2_2_27_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1epilogue_1_1warp_1_1VoltaTensorOpPolicy_3_01WarpShape___00_01gemm_1_1GemmShape_136ce744d4c1c6e8707f5a9785196194.html" target="_self">VoltaTensorOpPolicy&lt; WarpShape_, gemm::GemmShape&lt; 32, 32, 4 &gt;, float, layout::RowMajor &gt;</a></td><td class="desc">Partial specialization for row-major </td></tr>
<tr id="row_0_2_2_28_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1epilogue_1_1warp_1_1VoltaTensorOpPolicy_3_01WarpShape___00_01gemm_1_1GemmShape_1d48185f49e4d066f8e9327bf0856b7f.html" target="_self">VoltaTensorOpPolicy&lt; WarpShape_, gemm::GemmShape&lt; 32, 32, 4 &gt;, half_t, layout::RowMajor &gt;</a></td><td class="desc">Partial specialization for row-major </td></tr>
<tr id="row_0_2_3_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span id="arr_0_2_3_" class="arrow" onclick="toggleFolder('0_2_3_')">&#9658;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1epilogue_1_1EpilogueWorkspace.html" target="_self">EpilogueWorkspace</a></td><td class="desc"></td></tr>
<tr id="row_0_2_3_0_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1epilogue_1_1EpilogueWorkspace_1_1Params.html" target="_self">Params</a></td><td class="desc">Parameters structure </td></tr>
<tr id="row_0_2_3_1_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1epilogue_1_1EpilogueWorkspace_1_1SharedStorage.html" target="_self">SharedStorage</a></td><td class="desc">Shared storage allocation needed by the epilogue </td></tr>
<tr id="row_0_3_" style="display:none;"><td class="entry"><span style="width:16px;display:inline-block;">&#160;</span><span id="arr_0_3_" class="arrow" onclick="toggleFolder('0_3_')">&#9658;</span><span class="icona"><span class="icon">N</span></span><a class="el" href="namespacecutlass_1_1gemm.html" target="_self">gemm</a></td><td class="desc"></td></tr>
<tr id="row_0_3_0_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span id="arr_0_3_0_" class="arrow" onclick="toggleFolder('0_3_0_')">&#9658;</span><span class="icona"><span class="icon">N</span></span><a class="el" href="namespacecutlass_1_1gemm_1_1device.html" target="_self">device</a></td><td class="desc"></td></tr>
<tr id="row_0_3_0_0_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1gemm_1_1device_1_1DefaultGemmConfiguration.html" target="_self">DefaultGemmConfiguration</a></td><td class="desc"></td></tr>
<tr id="row_0_3_0_1_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1gemm_1_1device_1_1DefaultGemmConfiguration_3_01arch_1_1OpClassSimt_00_01ArchTag286687c5e6abe22d241f789fe344a465.html" target="_self">DefaultGemmConfiguration&lt; arch::OpClassSimt, ArchTag, ElementA, ElementB, ElementC, ElementAccumulator &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_3_0_2_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1gemm_1_1device_1_1DefaultGemmConfiguration_3_01arch_1_1OpClassSimt_00_01ArchTag3026e48abb8c905d1cc6d13d669700e4.html" target="_self">DefaultGemmConfiguration&lt; arch::OpClassSimt, ArchTag, int8_t, int8_t, ElementC, int32_t &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_3_0_3_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1gemm_1_1device_1_1DefaultGemmConfiguration_3_01arch_1_1OpClassTensorOp_00_01arc567cad318a31d04b70ea615d6321decd.html" target="_self">DefaultGemmConfiguration&lt; arch::OpClassTensorOp, arch::Sm70, ElementA, ElementB, ElementC, ElementAccumulator &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_3_0_4_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1gemm_1_1device_1_1DefaultGemmConfiguration_3_01arch_1_1OpClassTensorOp_00_01arcde61af9be1337dac1fdb210e7e7a6e01.html" target="_self">DefaultGemmConfiguration&lt; arch::OpClassTensorOp, arch::Sm75, ElementA, ElementB, ElementC, ElementAccumulator &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_3_0_5_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1gemm_1_1device_1_1DefaultGemmConfiguration_3_01arch_1_1OpClassTensorOp_00_01arc485a4f0b5a7d2d4ab2c1a24da6328048.html" target="_self">DefaultGemmConfiguration&lt; arch::OpClassTensorOp, arch::Sm75, int4b_t, int4b_t, ElementC, int32_t &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_3_0_6_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1gemm_1_1device_1_1DefaultGemmConfiguration_3_01arch_1_1OpClassTensorOp_00_01arc8e2604a56dff3a7595da9ee0604ae55e.html" target="_self">DefaultGemmConfiguration&lt; arch::OpClassTensorOp, arch::Sm75, int4b_t, uint4b_t, ElementC, int32_t &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_3_0_7_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1gemm_1_1device_1_1DefaultGemmConfiguration_3_01arch_1_1OpClassTensorOp_00_01arc4fada4957d463c80a2831e47f28157c4.html" target="_self">DefaultGemmConfiguration&lt; arch::OpClassTensorOp, arch::Sm75, int8_t, int8_t, ElementC, int32_t &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_3_0_8_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1gemm_1_1device_1_1DefaultGemmConfiguration_3_01arch_1_1OpClassTensorOp_00_01arc8ab5fd2693c6a6ec43e447acb07f784c.html" target="_self">DefaultGemmConfiguration&lt; arch::OpClassTensorOp, arch::Sm75, int8_t, uint8_t, ElementC, int32_t &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_3_0_9_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1gemm_1_1device_1_1DefaultGemmConfiguration_3_01arch_1_1OpClassTensorOp_00_01arcffcf31256aed23d4d8d0eab627bc0cad.html" target="_self">DefaultGemmConfiguration&lt; arch::OpClassTensorOp, arch::Sm75, uint4b_t, int4b_t, ElementC, int32_t &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_3_0_10_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1gemm_1_1device_1_1DefaultGemmConfiguration_3_01arch_1_1OpClassTensorOp_00_01arcb2e258b7bd321c633dd65d3ebcf6414a.html" target="_self">DefaultGemmConfiguration&lt; arch::OpClassTensorOp, arch::Sm75, uint4b_t, uint4b_t, ElementC, int32_t &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_3_0_11_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1gemm_1_1device_1_1DefaultGemmConfiguration_3_01arch_1_1OpClassTensorOp_00_01arcb27bf218007928652d5b803193eab473.html" target="_self">DefaultGemmConfiguration&lt; arch::OpClassTensorOp, arch::Sm75, uint8_t, int8_t, ElementC, int32_t &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_3_0_12_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1gemm_1_1device_1_1DefaultGemmConfiguration_3_01arch_1_1OpClassTensorOp_00_01arcfea0f3503156e8e3fba6456f0cedafdd.html" target="_self">DefaultGemmConfiguration&lt; arch::OpClassTensorOp, arch::Sm75, uint8_t, uint8_t, ElementC, int32_t &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_3_0_13_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1gemm_1_1device_1_1DefaultGemmConfiguration_3_01arch_1_1OpClassWmmaTensorOp_00_0884059ecad03bea3e86c4cf722226097.html" target="_self">DefaultGemmConfiguration&lt; arch::OpClassWmmaTensorOp, ArchTag, ElementA, ElementB, ElementC, ElementAccumulator &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_3_0_14_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span id="arr_0_3_0_14_" class="arrow" onclick="toggleFolder('0_3_0_14_')">&#9658;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1gemm_1_1device_1_1Gemm.html" target="_self">Gemm</a></td><td class="desc"></td></tr>
<tr id="row_0_3_0_14_0_" style="display:none;"><td class="entry"><span style="width:80px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1gemm_1_1device_1_1Gemm_1_1Arguments.html" target="_self">Arguments</a></td><td class="desc">Argument structure </td></tr>
<tr id="row_0_3_0_15_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span id="arr_0_3_0_15_" class="arrow" onclick="toggleFolder('0_3_0_15_')">&#9658;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1gemm_1_1device_1_1Gemm_3_01ElementA___00_01LayoutA___00_01ElementB___00_01Layout4d0960ae6b1d1bf19e6239dbd002249c.html" target="_self">Gemm&lt; ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, layout::ColumnMajor, ElementAccumulator_, OperatorClass_, ArchTag_, ThreadblockShape_, WarpShape_, InstructionShape_, EpilogueOutputOp_, ThreadblockSwizzle_, Stages, AlignmentA, AlignmentB, SplitKSerial, Operator_, IsBetaZero &gt;</a></td><td class="desc">Parital specialization for column-major output exchanges problem size and operand </td></tr>
<tr id="row_0_3_0_15_0_" style="display:none;"><td class="entry"><span style="width:80px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1gemm_1_1device_1_1Gemm_3_01ElementA___00_01LayoutA___00_01ElementB___00_01Layou1b211cc9c97c022d8fe10f2dd32c8709.html" target="_self">Arguments</a></td><td class="desc">Argument structure </td></tr>
<tr id="row_0_3_0_16_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span id="arr_0_3_0_16_" class="arrow" onclick="toggleFolder('0_3_0_16_')">&#9658;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched.html" target="_self">GemmBatched</a></td><td class="desc"></td></tr>
<tr id="row_0_3_0_16_0_" style="display:none;"><td class="entry"><span style="width:80px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1gemm_1_1device_1_1GemmBatched_1_1Arguments.html" target="_self">Arguments</a></td><td class="desc">Argument structure </td></tr>
<tr id="row_0_3_0_17_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span id="arr_0_3_0_17_" class="arrow" onclick="toggleFolder('0_3_0_17_')">&#9658;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_0c9bb6f4463ab6085e6008b5d5ad6abfd.html" target="_self">GemmBatched&lt; ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, layout::ColumnMajor, ElementAccumulator_, OperatorClass_, ArchTag_, ThreadblockShape_, WarpShape_, InstructionShape_, EpilogueOutputOp_, ThreadblockSwizzle_, Stages, AlignmentA, AlignmentB, Operator_ &gt;</a></td><td class="desc">Parital specialization for column-major output exchanges problem size and operand </td></tr>
<tr id="row_0_3_0_17_0_" style="display:none;"><td class="entry"><span style="width:80px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1gemm_1_1device_1_1GemmBatched_3_01ElementA___00_01LayoutA___00_01ElementB___00_213d78696663f4231cd52c6a277c60e5.html" target="_self">Arguments</a></td><td class="desc">Argument structure </td></tr>
<tr id="row_0_3_0_18_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span id="arr_0_3_0_18_" class="arrow" onclick="toggleFolder('0_3_0_18_')">&#9658;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1gemm_1_1device_1_1GemmComplex.html" target="_self">GemmComplex</a></td><td class="desc"></td></tr>
<tr id="row_0_3_0_18_0_" style="display:none;"><td class="entry"><span style="width:80px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1gemm_1_1device_1_1GemmComplex_1_1Arguments.html" target="_self">Arguments</a></td><td class="desc">Argument structure </td></tr>
<tr id="row_0_3_0_19_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span id="arr_0_3_0_19_" class="arrow" onclick="toggleFolder('0_3_0_19_')">&#9658;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1gemm_1_1device_1_1GemmComplex_3_01ElementA___00_01LayoutA___00_01ElementB___00_07c56401b4df75709ae636675d9980a9a.html" target="_self">GemmComplex&lt; ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, layout::ColumnMajor, ElementAccumulator_, OperatorClass_, ArchTag_, ThreadblockShape_, WarpShape_, InstructionShape_, EpilogueOutputOp_, ThreadblockSwizzle_, Stages, TransformA, TransformB, SplitKSerial &gt;</a></td><td class="desc">Parital specialization for column-major output exchanges problem size and operand </td></tr>
<tr id="row_0_3_0_19_0_" style="display:none;"><td class="entry"><span style="width:80px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1gemm_1_1device_1_1GemmComplex_3_01ElementA___00_01LayoutA___00_01ElementB___00_a3923967cafb5cb9774c320dc24baa77.html" target="_self">Arguments</a></td><td class="desc">Argument structure </td></tr>
<tr id="row_0_3_0_20_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span id="arr_0_3_0_20_" class="arrow" onclick="toggleFolder('0_3_0_20_')">&#9658;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1gemm_1_1device_1_1GemmSplitKParallel.html" target="_self">GemmSplitKParallel</a></td><td class="desc"></td></tr>
<tr id="row_0_3_0_20_0_" style="display:none;"><td class="entry"><span style="width:80px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1gemm_1_1device_1_1GemmSplitKParallel_1_1Arguments.html" target="_self">Arguments</a></td><td class="desc">Argument structure </td></tr>
<tr id="row_0_3_0_21_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span id="arr_0_3_0_21_" class="arrow" onclick="toggleFolder('0_3_0_21_')">&#9658;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1gemm_1_1device_1_1GemmSplitKParallel_3_01ElementA___00_01LayoutA___00_01ElementBbe7c1f7154ad5b5bf9d4d28301e2b457.html" target="_self">GemmSplitKParallel&lt; ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, layout::ColumnMajor, ElementAccumulator_, OperatorClass_, ArchTag_, ThreadblockShape_, WarpShape_, InstructionShape_, EpilogueOutputOp_, ConvertScaledOp_, ReductionOp_, ThreadblockSwizzle_, Stages, kAlignmentA, kAlignmentB, Operator_ &gt;</a></td><td class="desc">Partial specialization for column-major output </td></tr>
<tr id="row_0_3_0_21_0_" style="display:none;"><td class="entry"><span style="width:80px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1gemm_1_1device_1_1GemmSplitKParallel_3_01ElementA___00_01LayoutA___00_01Elementafcb1aeaf2035a7ac769d7acc233423b.html" target="_self">Arguments</a></td><td class="desc">Argument structure </td></tr>
<tr id="row_0_3_1_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span id="arr_0_3_1_" class="arrow" onclick="toggleFolder('0_3_1_')">&#9658;</span><span class="icona"><span class="icon">N</span></span><a class="el" href="namespacecutlass_1_1gemm_1_1kernel.html" target="_self">kernel</a></td><td class="desc"></td></tr>
<tr id="row_0_3_1_0_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span id="arr_0_3_1_0_" class="arrow" onclick="toggleFolder('0_3_1_0_')">&#9658;</span><span class="icona"><span class="icon">N</span></span><a class="el" href="namespacecutlass_1_1gemm_1_1kernel_1_1detail.html" target="_self">detail</a></td><td class="desc"></td></tr>
<tr id="row_0_3_1_0_0_" style="display:none;"><td class="entry"><span style="width:80px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1gemm_1_1kernel_1_1detail_1_1GemvBatchedStridedEpilogueScaling.html" target="_self">GemvBatchedStridedEpilogueScaling</a></td><td class="desc"></td></tr>
<tr id="row_0_3_1_1_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1gemm_1_1kernel_1_1DefaultGemm.html" target="_self">DefaultGemm</a></td><td class="desc"></td></tr>
<tr id="row_0_3_1_2_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1gemm_1_1kernel_1_1DefaultGemm_3_01ElementA_00_01layout_1_1ColumnMajorInterleave661fe54d13cc2c9153dcdf31e4beaa30.html" target="_self">DefaultGemm&lt; ElementA, layout::ColumnMajorInterleaved&lt; InterleavedK &gt;, kAlignmentA, ElementB, layout::RowMajorInterleaved&lt; InterleavedK &gt;, kAlignmentB, ElementC, layout::ColumnMajorInterleaved&lt; InterleavedK &gt;, int32_t, arch::OpClassTensorOp, arch::Sm75, ThreadblockShape, WarpShape, InstructionShape, EpilogueOutputOp, ThreadblockSwizzle, 2, SplitKSerial, Operator, IsBetaZero &gt;</a></td><td class="desc">Partial specialization for Turing Integer Matrix Multiply Interleaved layout </td></tr>
<tr id="row_0_3_1_3_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1gemm_1_1kernel_1_1DefaultGemm_3_01ElementA_00_01LayoutA_00_01kAlignmentA_00_01Edd80343e6570718ed237122e4ebf7fb5.html" target="_self">DefaultGemm&lt; ElementA, LayoutA, kAlignmentA, ElementB, LayoutB, kAlignmentB, ElementC, layout::RowMajor, ElementAccumulator, arch::OpClassSimt, ArchTag, ThreadblockShape, WarpShape, GemmShape&lt; 1, 1, 1 &gt;, EpilogueOutputOp, ThreadblockSwizzle, 2, SplitKSerial, Operator &gt;</a></td><td class="desc">Partial specialization for SIMT </td></tr>
<tr id="row_0_3_1_4_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1gemm_1_1kernel_1_1DefaultGemm_3_01ElementA_00_01LayoutA_00_01kAlignmentA_00_01E044b039b2fe402f29b04a9f5feee5342.html" target="_self">DefaultGemm&lt; ElementA, LayoutA, kAlignmentA, ElementB, LayoutB, kAlignmentB, ElementC, layout::RowMajor, ElementAccumulator, arch::OpClassTensorOp, arch::Sm70, ThreadblockShape, WarpShape, GemmShape&lt; 8, 8, 4 &gt;, EpilogueOutputOp, ThreadblockSwizzle, 2, SplitKSerial, Operator &gt;</a></td><td class="desc">Partial specialization for Volta architecture </td></tr>
<tr id="row_0_3_1_5_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1gemm_1_1kernel_1_1DefaultGemm_3_01ElementA_00_01LayoutA_00_01kAlignmentA_00_01E5d78d37a9ae2ec08d7d477d571df036e.html" target="_self">DefaultGemm&lt; ElementA, LayoutA, kAlignmentA, ElementB, LayoutB, kAlignmentB, ElementC, layout::RowMajor, ElementAccumulator, arch::OpClassTensorOp, arch::Sm75, ThreadblockShape, WarpShape, InstructionShape, EpilogueOutputOp, ThreadblockSwizzle, 2, SplitKSerial, Operator &gt;</a></td><td class="desc">Partial specialization for Turing Architecture </td></tr>
<tr id="row_0_3_1_6_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1gemm_1_1kernel_1_1DefaultGemm_3_01int8__t_00_01LayoutA_00_01kAlignmentA_00_01inf48440732c1c5f42ddbfaba179861815.html" target="_self">DefaultGemm&lt; int8_t, LayoutA, kAlignmentA, int8_t, LayoutB, kAlignmentB, ElementC, LayoutC, ElementAccumulator, arch::OpClassSimt, ArchTag, ThreadblockShape, WarpShape, GemmShape&lt; 1, 1, 4 &gt;, EpilogueOutputOp, ThreadblockSwizzle, 2, SplitKSerial, Operator, false &gt;</a></td><td class="desc">Partial specialization for SIMT DP4A </td></tr>
<tr id="row_0_3_1_7_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1gemm_1_1kernel_1_1DefaultGemmSplitKParallel.html" target="_self">DefaultGemmSplitKParallel</a></td><td class="desc"></td></tr>
<tr id="row_0_3_1_8_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1gemm_1_1kernel_1_1DefaultGemv.html" target="_self">DefaultGemv</a></td><td class="desc"></td></tr>
<tr id="row_0_3_1_9_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span id="arr_0_3_1_9_" class="arrow" onclick="toggleFolder('0_3_1_9_')">&#9658;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1gemm_1_1kernel_1_1Gemm.html" target="_self">Gemm</a></td><td class="desc"></td></tr>
<tr id="row_0_3_1_9_0_" style="display:none;"><td class="entry"><span style="width:80px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1gemm_1_1kernel_1_1Gemm_1_1Params.html" target="_self">Params</a></td><td class="desc">Parameters structure </td></tr>
<tr id="row_0_3_1_9_1_" style="display:none;"><td class="entry"><span style="width:80px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="unioncutlass_1_1gemm_1_1kernel_1_1Gemm_1_1SharedStorage.html" target="_self">SharedStorage</a></td><td class="desc">Shared memory storage structure </td></tr>
<tr id="row_0_3_1_10_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span id="arr_0_3_1_10_" class="arrow" onclick="toggleFolder('0_3_1_10_')">&#9658;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1gemm_1_1kernel_1_1GemmBatched.html" target="_self">GemmBatched</a></td><td class="desc"></td></tr>
<tr id="row_0_3_1_10_0_" style="display:none;"><td class="entry"><span style="width:80px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1gemm_1_1kernel_1_1GemmBatched_1_1Params.html" target="_self">Params</a></td><td class="desc">Parameters structure </td></tr>
<tr id="row_0_3_1_10_1_" style="display:none;"><td class="entry"><span style="width:80px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="unioncutlass_1_1gemm_1_1kernel_1_1GemmBatched_1_1SharedStorage.html" target="_self">SharedStorage</a></td><td class="desc">Shared memory storage structure </td></tr>
<tr id="row_0_3_1_11_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span id="arr_0_3_1_11_" class="arrow" onclick="toggleFolder('0_3_1_11_')">&#9658;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1gemm_1_1kernel_1_1GemmSplitKParallel.html" target="_self">GemmSplitKParallel</a></td><td class="desc"></td></tr>
<tr id="row_0_3_1_11_0_" style="display:none;"><td class="entry"><span style="width:80px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1gemm_1_1kernel_1_1GemmSplitKParallel_1_1Params.html" target="_self">Params</a></td><td class="desc">Parameters structure </td></tr>
<tr id="row_0_3_1_11_1_" style="display:none;"><td class="entry"><span style="width:80px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="unioncutlass_1_1gemm_1_1kernel_1_1GemmSplitKParallel_1_1SharedStorage.html" target="_self">SharedStorage</a></td><td class="desc">Shared memory storage structure </td></tr>
<tr id="row_0_3_2_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span id="arr_0_3_2_" class="arrow" onclick="toggleFolder('0_3_2_')">&#9658;</span><span class="icona"><span class="icon">N</span></span><a class="el" href="namespacecutlass_1_1gemm_1_1thread.html" target="_self">thread</a></td><td class="desc"></td></tr>
<tr id="row_0_3_2_0_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span id="arr_0_3_2_0_" class="arrow" onclick="toggleFolder('0_3_2_0_')">&#9658;</span><span class="icona"><span class="icon">N</span></span><a class="el" href="namespacecutlass_1_1gemm_1_1thread_1_1detail.html" target="_self">detail</a></td><td class="desc"></td></tr>
<tr id="row_0_3_2_0_0_" style="display:none;"><td class="entry"><span style="width:80px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1gemm_1_1thread_1_1detail_1_1EnableMma__Crow__SM60.html" target="_self">EnableMma_Crow_SM60</a></td><td class="desc">Determines whether to enable thread::Gemm&lt;&gt; specializations compatible with SM50 </td></tr>
<tr id="row_0_3_2_0_1_" style="display:none;"><td class="entry"><span style="width:80px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1gemm_1_1thread_1_1detail_1_1Mma__HFMA2.html" target="_self">Mma_HFMA2</a></td><td class="desc">Structure to compute the matrix product for HFMA </td></tr>
<tr id="row_0_3_2_0_2_" style="display:none;"><td class="entry"><span style="width:80px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1gemm_1_1thread_1_1detail_1_1Mma__HFMA2_3_01Shape_00_01layout_1_1ColumnMajor_00_72621f7ab9ae4a4ba4fe9725cf8e89c1.html" target="_self">Mma_HFMA2&lt; Shape, layout::ColumnMajor, layout::ColumnMajor, layout::ColumnMajor, true &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_3_2_0_3_" style="display:none;"><td class="entry"><span style="width:80px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1gemm_1_1thread_1_1detail_1_1Mma__HFMA2_3_01Shape_00_01layout_1_1ColumnMajor_00_94c813e3bbfb6f9857c155166f772687.html" target="_self">Mma_HFMA2&lt; Shape, layout::ColumnMajor, layout::ColumnMajor, layout::RowMajor, true &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_3_2_0_4_" style="display:none;"><td class="entry"><span style="width:80px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1gemm_1_1thread_1_1detail_1_1Mma__HFMA2_3_01Shape_00_01layout_1_1ColumnMajor_00_17070298bc4cced0a1b98aee2bb6b455.html" target="_self">Mma_HFMA2&lt; Shape, layout::ColumnMajor, layout::RowMajor, layout::ColumnMajor, true &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_3_2_0_5_" style="display:none;"><td class="entry"><span style="width:80px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1gemm_1_1thread_1_1detail_1_1Mma__HFMA2_3_01Shape_00_01layout_1_1ColumnMajor_00_bf6d29bb09a025e7b96942809743e28a.html" target="_self">Mma_HFMA2&lt; Shape, layout::ColumnMajor, layout::RowMajor, layout::RowMajor, true &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_3_2_0_6_" style="display:none;"><td class="entry"><span style="width:80px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1gemm_1_1thread_1_1detail_1_1Mma__HFMA2_3_01Shape_00_01layout_1_1RowMajor_00_01l26a133b13650c1d058273e3649f60f04.html" target="_self">Mma_HFMA2&lt; Shape, layout::RowMajor, layout::ColumnMajor, layout::ColumnMajor, true &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_3_2_0_7_" style="display:none;"><td class="entry"><span style="width:80px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1gemm_1_1thread_1_1detail_1_1Mma__HFMA2_3_01Shape_00_01layout_1_1RowMajor_00_01lbba3a796be96a0276693ef6b259ecc4a.html" target="_self">Mma_HFMA2&lt; Shape, layout::RowMajor, layout::ColumnMajor, layout::RowMajor, true &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_3_2_0_8_" style="display:none;"><td class="entry"><span style="width:80px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1gemm_1_1thread_1_1detail_1_1Mma__HFMA2_3_01Shape_00_01layout_1_1RowMajor_00_01l2aa4d2fd2e940e0d0cf7c47bc8f6017c.html" target="_self">Mma_HFMA2&lt; Shape, layout::RowMajor, layout::RowMajor, layout::ColumnMajor, true &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_3_2_0_9_" style="display:none;"><td class="entry"><span style="width:80px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1gemm_1_1thread_1_1detail_1_1Mma__HFMA2_3_01Shape_00_01layout_1_1RowMajor_00_01l086c058a15d6c79558e4f3d9ff1dc148.html" target="_self">Mma_HFMA2&lt; Shape, layout::RowMajor, layout::RowMajor, layout::RowMajor, true &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_3_2_0_10_" style="display:none;"><td class="entry"><span style="width:80px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1gemm_1_1thread_1_1detail_1_1Mma__HFMA2_3_01Shape_00_01LayoutA_00_01LayoutB_00_0e1104c65871c539155bd3a0c7631928b.html" target="_self">Mma_HFMA2&lt; Shape, LayoutA, LayoutB, layout::ColumnMajor, false &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_3_2_0_11_" style="display:none;"><td class="entry"><span style="width:80px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1gemm_1_1thread_1_1detail_1_1Mma__HFMA2_3_01Shape_00_01LayoutA_00_01LayoutB_00_07ac147cb320ee0d28ff8e78eb4cd330e.html" target="_self">Mma_HFMA2&lt; Shape, LayoutA, LayoutB, layout::RowMajor, false &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_3_2_1_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1gemm_1_1thread_1_1Mma.html" target="_self">Mma</a></td><td class="desc">Structure to compute the matrix product </td></tr>
<tr id="row_0_3_2_2_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1gemm_1_1thread_1_1Mma_3_01Shape___00_01ElementA___00_01LayoutA___00_01ElementB_e41c1cd6078b6d1347fac239b0639d56.html" target="_self">Mma&lt; Shape_, ElementA_, LayoutA_, ElementB_, LayoutB_, ElementC_, LayoutC_, arch::OpMultiplyAdd, bool &gt;</a></td><td class="desc">Gemplate that handles conventional layouts for FFMA and DFMA GEMM </td></tr>
<tr id="row_0_3_2_3_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1gemm_1_1thread_1_1Mma_3_01Shape___00_01half__t_00_01LayoutA_00_01half__t_00_01L066c9d2371712cdf0cac099ca9bcc578.html" target="_self">Mma&lt; Shape_, half_t, LayoutA, half_t, LayoutB, half_t, LayoutC, arch::OpMultiplyAdd &gt;</a></td><td class="desc">Structure to compute the matrix product </td></tr>
<tr id="row_0_3_2_4_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1gemm_1_1thread_1_1Mma_3_01Shape___00_01half__t_00_01LayoutA___00_01half__t_00_088f0e99e501b6012297eb30b4e89bcea.html" target="_self">Mma&lt; Shape_, half_t, LayoutA_, half_t, LayoutB_, half_t, layout::RowMajor, arch::OpMultiplyAdd, typename platform::enable_if&lt; detail::EnableMma_Crow_SM60&lt; LayoutA_, LayoutB_ &gt;::value &gt;::type &gt;</a></td><td class="desc">Computes matrix product when C is row-major </td></tr>
<tr id="row_0_3_2_5_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1gemm_1_1thread_1_1Mma_3_01Shape___00_01int8__t_00_01layout_1_1ColumnMajor_00_013f3785e722edc6e9aab6f866309b8623.html" target="_self">Mma&lt; Shape_, int8_t, layout::ColumnMajor, int8_t, layout::RowMajor, int32_t, LayoutC_, arch::OpMultiplyAdd, int8_t &gt;</a></td><td class="desc">Gemplate that handles conventional layouts for IDP4A </td></tr>
<tr id="row_0_3_2_6_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1gemm_1_1thread_1_1Mma_3_01Shape___00_01int8__t_00_01layout_1_1RowMajor_00_01int89c659e7faf47264972bdba6cd80f42b.html" target="_self">Mma&lt; Shape_, int8_t, layout::RowMajor, int8_t, layout::ColumnMajor, int32_t, LayoutC_, arch::OpMultiplyAdd, bool &gt;</a></td><td class="desc">Gemplate that handles conventional layouts for IDP4A </td></tr>
<tr id="row_0_3_2_7_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1gemm_1_1thread_1_1MmaGeneric.html" target="_self">MmaGeneric</a></td><td class="desc">Gemplate that handles all packed matrix layouts </td></tr>
<tr id="row_0_3_3_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span id="arr_0_3_3_" class="arrow" onclick="toggleFolder('0_3_3_')">&#9658;</span><span class="icona"><span class="icon">N</span></span><a class="el" href="namespacecutlass_1_1gemm_1_1threadblock.html" target="_self">threadblock</a></td><td class="desc"></td></tr>
<tr id="row_0_3_3_0_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1gemm_1_1threadblock_1_1DefaultGemvCore.html" target="_self">DefaultGemvCore</a></td><td class="desc"></td></tr>
<tr id="row_0_3_3_1_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1gemm_1_1threadblock_1_1DefaultMma.html" target="_self">DefaultMma</a></td><td class="desc"></td></tr>
<tr id="row_0_3_3_2_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1gemm_1_1threadblock_1_1DefaultMma_3_01ElementA_00_01LayoutA_00_01kAlignmentA_0010764e1fd5a3251a57eddafbd83eab8e.html" target="_self">DefaultMma&lt; ElementA, LayoutA, kAlignmentA, ElementB, LayoutB, kAlignmentB, ElementAccumulator, layout::ColumnMajorInterleaved&lt; InterleavedK &gt;, OperatorClass, ArchTag, ThreadblockShape, WarpShape, InstructionShape, 2, Operator, true &gt;</a></td><td class="desc">Specialization for column-major-interleaved output </td></tr>
<tr id="row_0_3_3_3_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1gemm_1_1threadblock_1_1DefaultMma_3_01ElementA_00_01LayoutA_00_01kAlignmentA_00c67c16f9881e4f2fda76d8ed83ebabd6.html" target="_self">DefaultMma&lt; ElementA, LayoutA, kAlignmentA, ElementB, LayoutB, kAlignmentB, ElementAccumulator, layout::RowMajor, arch::OpClassSimt, ArchTag, ThreadblockShape, WarpShape, InstructionShape, 2, Operator, false &gt;</a></td><td class="desc">Specialization for row-major output (OperatorClass Simt) </td></tr>
<tr id="row_0_3_3_4_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1gemm_1_1threadblock_1_1DefaultMma_3_01ElementA_00_01LayoutA_00_01kAlignmentA_00ce36642cae579bce6605ff8edde3c6ab.html" target="_self">DefaultMma&lt; ElementA, LayoutA, kAlignmentA, ElementB, LayoutB, kAlignmentB, ElementAccumulator, layout::RowMajor, arch::OpClassTensorOp, ArchTag, ThreadblockShape, WarpShape, InstructionShape, 2, Operator, false &gt;</a></td><td class="desc">Specialization for row-major output (OperatorClass Simt) </td></tr>
<tr id="row_0_3_3_5_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1gemm_1_1threadblock_1_1DefaultMma_3_01int8__t_00_01LayoutA_00_01kAlignmentA_00_07e7230d4011ada5e22cfcb29103b696.html" target="_self">DefaultMma&lt; int8_t, LayoutA, kAlignmentA, int8_t, LayoutB, kAlignmentB, ElementAccumulator, layout::RowMajor, arch::OpClassSimt, ArchTag, ThreadblockShape, WarpShape, GemmShape&lt; 1, 1, 4 &gt;, 2, Operator, false &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_3_3_6_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1gemm_1_1threadblock_1_1DefaultMmaCore.html" target="_self">DefaultMmaCore</a></td><td class="desc"></td></tr>
<tr id="row_0_3_3_7_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1gemm_1_1threadblock_1_1DefaultMmaCore_3_01Shape___00_01WarpShape___00_01GemmShab94a11a77dd0565102710907089acee0.html" target="_self">DefaultMmaCore&lt; Shape_, WarpShape_, GemmShape&lt; 1, 1, 1 &gt;, ElementA_, layout::ColumnMajor, ElementB_, layout::ColumnMajor, ElementC_, LayoutC_, arch::OpClassSimt, 2, Operator_ &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_3_3_8_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1gemm_1_1threadblock_1_1DefaultMmaCore_3_01Shape___00_01WarpShape___00_01GemmShafafd5c61db86cbfe90863578ddd11092.html" target="_self">DefaultMmaCore&lt; Shape_, WarpShape_, GemmShape&lt; 1, 1, 1 &gt;, ElementA_, layout::ColumnMajor, ElementB_, layout::RowMajor, ElementC_, LayoutC_, arch::OpClassSimt, 2, Operator_ &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_3_3_9_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1gemm_1_1threadblock_1_1DefaultMmaCore_3_01Shape___00_01WarpShape___00_01GemmSha46446d1e3871e31d2e728f710d78c8c1.html" target="_self">DefaultMmaCore&lt; Shape_, WarpShape_, GemmShape&lt; 1, 1, 1 &gt;, ElementA_, layout::ColumnMajor, ElementB_, layout::RowMajor, ElementC_, LayoutC_, arch::OpClassSimt, 2, Operator_, &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_3_3_10_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1gemm_1_1threadblock_1_1DefaultMmaCore_3_01Shape___00_01WarpShape___00_01GemmSha8da7a0cfbbe859b701fdd9f2b8566aa7.html" target="_self">DefaultMmaCore&lt; Shape_, WarpShape_, GemmShape&lt; 1, 1, 1 &gt;, ElementA_, layout::RowMajor, ElementB_, layout::ColumnMajor, ElementC_, LayoutC_, arch::OpClassSimt, 2, Operator_ &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_3_3_11_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1gemm_1_1threadblock_1_1DefaultMmaCore_3_01Shape___00_01WarpShape___00_01GemmSha84e9f8afb6a4ca9f5dcd219b182d16e7.html" target="_self">DefaultMmaCore&lt; Shape_, WarpShape_, GemmShape&lt; 1, 1, 1 &gt;, ElementA_, layout::RowMajor, ElementB_, layout::RowMajor, ElementC_, LayoutC_, arch::OpClassSimt, 2, Operator_ &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_3_3_12_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1gemm_1_1threadblock_1_1DefaultMmaCore_3_01Shape___00_01WarpShape___00_01GemmSha2c0d0b7cdb5c4bcb11e83c058eb65345.html" target="_self">DefaultMmaCore&lt; Shape_, WarpShape_, GemmShape&lt; 1, 1, 4 &gt;, int8_t, layout::ColumnMajor, int8_t, layout::ColumnMajor, ElementC_, LayoutC_, arch::OpClassSimt, 2, Operator_ &gt;</a></td><td class="desc">Partial specialization: </td></tr>
<tr id="row_0_3_3_13_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1gemm_1_1threadblock_1_1DefaultMmaCore_3_01Shape___00_01WarpShape___00_01GemmSha34a52cc7b2942e8c290f0032b6779b52.html" target="_self">DefaultMmaCore&lt; Shape_, WarpShape_, GemmShape&lt; 1, 1, 4 &gt;, int8_t, layout::ColumnMajor, int8_t, layout::RowMajor, ElementC_, LayoutC_, arch::OpClassSimt, 2, Operator_ &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_3_3_14_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1gemm_1_1threadblock_1_1DefaultMmaCore_3_01Shape___00_01WarpShape___00_01GemmShaaf312aafe9da92ea9d417bcc12a8e7dc.html" target="_self">DefaultMmaCore&lt; Shape_, WarpShape_, GemmShape&lt; 1, 1, 4 &gt;, int8_t, layout::RowMajor, int8_t, layout::ColumnMajor, ElementC_, LayoutC_, arch::OpClassSimt, 2, Operator_ &gt;</a></td><td class="desc">Partial specialization: </td></tr>
<tr id="row_0_3_3_15_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1gemm_1_1threadblock_1_1DefaultMmaCore_3_01Shape___00_01WarpShape___00_01GemmSha863d4139ccaa713bc4bde32c425f4067.html" target="_self">DefaultMmaCore&lt; Shape_, WarpShape_, GemmShape&lt; 1, 1, 4 &gt;, int8_t, layout::RowMajor, int8_t, layout::RowMajor, ElementC_, LayoutC_, arch::OpClassSimt, 2, Operator_ &gt;</a></td><td class="desc">Partial specialization: </td></tr>
<tr id="row_0_3_3_16_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1gemm_1_1threadblock_1_1DefaultMmaCore_3_01Shape___00_01WarpShape___00_01GemmShaf03a122202ad10acdc96f280106d678b.html" target="_self">DefaultMmaCore&lt; Shape_, WarpShape_, GemmShape&lt; 8, 8, 4 &gt;, ElementA_, layout::ColumnMajor, ElementB_, layout::ColumnMajor, ElementC_, LayoutC_, arch::OpClassTensorOp, 2, Operator_ &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_3_3_17_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1gemm_1_1threadblock_1_1DefaultMmaCore_3_01Shape___00_01WarpShape___00_01GemmSha69bef08ea63dd930f99d9788105873dd.html" target="_self">DefaultMmaCore&lt; Shape_, WarpShape_, GemmShape&lt; 8, 8, 4 &gt;, ElementA_, layout::ColumnMajor, ElementB_, layout::RowMajor, ElementC_, LayoutC_, arch::OpClassTensorOp, 2, Operator_ &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_3_3_18_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1gemm_1_1threadblock_1_1DefaultMmaCore_3_01Shape___00_01WarpShape___00_01GemmSha3adf608332a8c9ee7014fced0da8a9ca.html" target="_self">DefaultMmaCore&lt; Shape_, WarpShape_, GemmShape&lt; 8, 8, 4 &gt;, ElementA_, layout::RowMajor, ElementB_, layout::ColumnMajor, ElementC_, LayoutC_, arch::OpClassTensorOp, 2, Operator_ &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_3_3_19_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1gemm_1_1threadblock_1_1DefaultMmaCore_3_01Shape___00_01WarpShape___00_01GemmShab7edfba3cdf43a07e3c4d719d87565a4.html" target="_self">DefaultMmaCore&lt; Shape_, WarpShape_, GemmShape&lt; 8, 8, 4 &gt;, ElementA_, layout::RowMajor, ElementB_, layout::RowMajor, ElementC_, LayoutC_, arch::OpClassTensorOp, 2, Operator_ &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_3_3_20_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1gemm_1_1threadblock_1_1DefaultMmaCore_3_01Shape___00_01WarpShape___00_01Instruc803d38bc1e4618c07c47f54c87ae2678.html" target="_self">DefaultMmaCore&lt; Shape_, WarpShape_, InstructionShape_, ElementA_, layout::ColumnMajor, ElementB_, layout::ColumnMajor, ElementC_, LayoutC_, arch::OpClassTensorOp, 2, Operator_ &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_3_3_21_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1gemm_1_1threadblock_1_1DefaultMmaCore_3_01Shape___00_01WarpShape___00_01Instrucf60fe02fcdd80d28b7fd419133465dcc.html" target="_self">DefaultMmaCore&lt; Shape_, WarpShape_, InstructionShape_, ElementA_, layout::ColumnMajor, ElementB_, layout::RowMajor, ElementC_, LayoutC_, arch::OpClassTensorOp, 2, Operator_ &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_3_3_22_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1gemm_1_1threadblock_1_1DefaultMmaCore_3_01Shape___00_01WarpShape___00_01Instruc2bf00737f4ad0a9da9a8be6d3e66c152.html" target="_self">DefaultMmaCore&lt; Shape_, WarpShape_, InstructionShape_, ElementA_, layout::ColumnMajorInterleaved&lt; InterleavedK &gt;, ElementB_, layout::RowMajorInterleaved&lt; InterleavedK &gt;, ElementC_, LayoutC_, arch::OpClassTensorOp, 2, Operator_, AccumulatorsInRowMajor &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_3_3_23_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1gemm_1_1threadblock_1_1DefaultMmaCore_3_01Shape___00_01WarpShape___00_01Instruc24092ddc01fc83dabb7db4c14880fe60.html" target="_self">DefaultMmaCore&lt; Shape_, WarpShape_, InstructionShape_, ElementA_, layout::RowMajor, ElementB_, layout::ColumnMajor, ElementC_, LayoutC_, arch::OpClassTensorOp, 2, Operator_ &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_3_3_24_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1gemm_1_1threadblock_1_1DefaultMmaCore_3_01Shape___00_01WarpShape___00_01Instruc4fee9f2965b8468bfb42b94a74527d22.html" target="_self">DefaultMmaCore&lt; Shape_, WarpShape_, InstructionShape_, ElementA_, layout::RowMajor, ElementB_, layout::RowMajor, ElementC_, LayoutC_, arch::OpClassTensorOp, 2, Operator_ &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_3_3_25_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1gemm_1_1threadblock_1_1GemmBatchedIdentityThreadblockSwizzle.html" target="_self">GemmBatchedIdentityThreadblockSwizzle</a></td><td class="desc">Threadblock swizzling function for batched GEMMs </td></tr>
<tr id="row_0_3_3_26_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1gemm_1_1threadblock_1_1GemmHorizontalThreadblockSwizzle.html" target="_self">GemmHorizontalThreadblockSwizzle</a></td><td class="desc">Threadblock swizzling function for GEMMs </td></tr>
<tr id="row_0_3_3_27_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1gemm_1_1threadblock_1_1GemmIdentityThreadblockSwizzle.html" target="_self">GemmIdentityThreadblockSwizzle</a></td><td class="desc">Threadblock swizzling function for GEMMs </td></tr>
<tr id="row_0_3_3_28_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1gemm_1_1threadblock_1_1GemmSplitKHorizontalThreadblockSwizzle.html" target="_self">GemmSplitKHorizontalThreadblockSwizzle</a></td><td class="desc">Threadblock swizzling function for split-K GEMMs </td></tr>
<tr id="row_0_3_3_29_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1gemm_1_1threadblock_1_1GemmSplitKIdentityThreadblockSwizzle.html" target="_self">GemmSplitKIdentityThreadblockSwizzle</a></td><td class="desc">Threadblock swizzling function for split-K GEMMs </td></tr>
<tr id="row_0_3_3_30_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1gemm_1_1threadblock_1_1Gemv.html" target="_self">Gemv</a></td><td class="desc">Structure to compute the matrix-vector product using SIMT math instructions </td></tr>
<tr id="row_0_3_3_31_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1gemm_1_1threadblock_1_1GemvBatchedStridedThreadblockDefaultSwizzle.html" target="_self">GemvBatchedStridedThreadblockDefaultSwizzle</a></td><td class="desc">Threadblock swizzling function for batched GEMVs </td></tr>
<tr id="row_0_3_3_32_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span id="arr_0_3_3_32_" class="arrow" onclick="toggleFolder('0_3_3_32_')">&#9658;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1gemm_1_1threadblock_1_1MmaBase.html" target="_self">MmaBase</a></td><td class="desc"></td></tr>
<tr id="row_0_3_3_32_0_" style="display:none;"><td class="entry"><span style="width:80px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1gemm_1_1threadblock_1_1MmaBase_1_1SharedStorage.html" target="_self">SharedStorage</a></td><td class="desc">Shared storage object needed by threadblock-scoped GEMM </td></tr>
<tr id="row_0_3_3_33_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1gemm_1_1threadblock_1_1MmaPipelined.html" target="_self">MmaPipelined</a></td><td class="desc">Structure to compute the matrix product targeting CUDA cores and SIMT math instructions </td></tr>
<tr id="row_0_3_3_34_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1gemm_1_1threadblock_1_1MmaPolicy.html" target="_self">MmaPolicy</a></td><td class="desc">Policy object describing MmaTensorOp </td></tr>
<tr id="row_0_3_3_35_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1gemm_1_1threadblock_1_1MmaSingleStage.html" target="_self">MmaSingleStage</a></td><td class="desc">Structure to compute the matrix product targeting CUDA cores and SIMT math instructions </td></tr>
<tr id="row_0_3_4_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span id="arr_0_3_4_" class="arrow" onclick="toggleFolder('0_3_4_')">&#9658;</span><span class="icona"><span class="icon">N</span></span><a class="el" href="namespacecutlass_1_1gemm_1_1warp.html" target="_self">warp</a></td><td class="desc"></td></tr>
<tr id="row_0_3_4_0_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1gemm_1_1warp_1_1DefaultMmaTensorOp.html" target="_self">DefaultMmaTensorOp</a></td><td class="desc">Partial specialization for m-by-n-by-kgroup </td></tr>
<tr id="row_0_3_4_1_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1gemm_1_1warp_1_1MmaComplexTensorOp.html" target="_self">MmaComplexTensorOp</a></td><td class="desc"></td></tr>
<tr id="row_0_3_4_2_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1gemm_1_1warp_1_1MmaComplexTensorOp_3_01Shape___00_01complex_3_01RealElementA_01_146441010dad1f40eb51b6dae3ded216.html" target="_self">MmaComplexTensorOp&lt; Shape_, complex&lt; RealElementA &gt;, LayoutA_, complex&lt; RealElementB &gt;, LayoutB_, complex&lt; RealElementC &gt;, LayoutC_, Policy_, TransformA, TransformB, Enable &gt;</a></td><td class="desc">Partial specialization for complex*complex+complex =&gt; complex using real-valued TensorOps </td></tr>
<tr id="row_0_3_4_3_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1gemm_1_1warp_1_1MmaSimt.html" target="_self">MmaSimt</a></td><td class="desc">Structure to compute the matrix product targeting CUDA cores and SIMT math instructions </td></tr>
<tr id="row_0_3_4_4_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1gemm_1_1warp_1_1MmaSimtPolicy.html" target="_self">MmaSimtPolicy</a></td><td class="desc">Describes the arrangement and configuration of per-lane operations in warp-level matrix multiply </td></tr>
<tr id="row_0_3_4_5_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1gemm_1_1warp_1_1MmaSimtTileIterator.html" target="_self">MmaSimtTileIterator</a></td><td class="desc"></td></tr>
<tr id="row_0_3_4_6_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1gemm_1_1warp_1_1MmaSimtTileIterator_3_01Shape___00_01Operand_1_1kA_00_01Element_67ca7e11a38e38f2c51b84767654a90f.html" target="_self">MmaSimtTileIterator&lt; Shape_, Operand::kA, Element_, layout::ColumnMajor, Policy_, PartitionsK, PartitionGroupSize &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_3_4_7_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1gemm_1_1warp_1_1MmaSimtTileIterator_3_01Shape___00_01Operand_1_1kA_00_01Element_f0ce904a9294556f15e1cc9cf7c99a93.html" target="_self">MmaSimtTileIterator&lt; Shape_, Operand::kA, Element_, layout::ColumnMajorInterleaved&lt; 4 &gt;, Policy_, PartitionsK, PartitionGroupSize &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_3_4_8_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1gemm_1_1warp_1_1MmaSimtTileIterator_3_01Shape___00_01Operand_1_1kB_00_01Element_ea0a4e7ce3cd5d25cabf79383efdf4d9.html" target="_self">MmaSimtTileIterator&lt; Shape_, Operand::kB, Element_, layout::RowMajor, Policy_, PartitionsK, PartitionGroupSize &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_3_4_9_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1gemm_1_1warp_1_1MmaSimtTileIterator_3_01Shape___00_01Operand_1_1kB_00_01Element_ada156b62fcbdce47009c5bf1321c92c.html" target="_self">MmaSimtTileIterator&lt; Shape_, Operand::kB, Element_, layout::RowMajorInterleaved&lt; 4 &gt;, Policy_, PartitionsK, PartitionGroupSize &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_3_4_10_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1gemm_1_1warp_1_1MmaSimtTileIterator_3_01Shape___00_01Operand_1_1kC_00_01Element_4ccafbc821b3a55cd532602442a74031.html" target="_self">MmaSimtTileIterator&lt; Shape_, Operand::kC, Element_, layout::ColumnMajor, Policy_ &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_3_4_11_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1gemm_1_1warp_1_1MmaSimtTileIterator_3_01Shape___00_01Operand_1_1kC_00_01Element_8f92ea79e85febb67169c4b2d94b1b20.html" target="_self">MmaSimtTileIterator&lt; Shape_, Operand::kC, Element_, layout::RowMajor, Policy_ &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_3_4_12_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1gemm_1_1warp_1_1MmaTensorOp.html" target="_self">MmaTensorOp</a></td><td class="desc">Structure to compute the matrix product targeting CUDA cores and SIMT math instructions </td></tr>
<tr id="row_0_3_4_13_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1gemm_1_1warp_1_1MmaTensorOpAccumulatorTileIterator.html" target="_self">MmaTensorOpAccumulatorTileIterator</a></td><td class="desc"></td></tr>
<tr id="row_0_3_4_14_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span id="arr_0_3_4_14_" class="arrow" onclick="toggleFolder('0_3_4_14_')">&#9658;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1gemm_1_1warp_1_1MmaTensorOpAccumulatorTileIterator_3_01Shape___00_01Element___008f607b871a2b3d854eb4def64712c042.html" target="_self">MmaTensorOpAccumulatorTileIterator&lt; Shape_, Element_, cutlass::layout::ColumnMajor, InstructionShape_, OpDelta_ &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_3_4_14_0_" style="display:none;"><td class="entry"><span style="width:80px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1gemm_1_1warp_1_1MmaTensorOpAccumulatorTileIterator_3_01Shape___00_01Element___0d35fa5dc4e4b4f72784c943fd857fc1d.html" target="_self">Policy</a></td><td class="desc">Internal structure of iterator - made public to enable introspection </td></tr>
<tr id="row_0_3_4_15_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span id="arr_0_3_4_15_" class="arrow" onclick="toggleFolder('0_3_4_15_')">&#9658;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1gemm_1_1warp_1_1MmaTensorOpAccumulatorTileIterator_3_01Shape___00_01Element___00027dabdc144edd6276f664ca74088510.html" target="_self">MmaTensorOpAccumulatorTileIterator&lt; Shape_, Element_, cutlass::layout::ColumnMajorInterleaved&lt; InterleavedN &gt;, InstructionShape_, OpDelta_ &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_3_4_15_0_" style="display:none;"><td class="entry"><span style="width:80px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1gemm_1_1warp_1_1MmaTensorOpAccumulatorTileIterator_3_01Shape___00_01Element___03822d9be37f3725022005a5434441f22.html" target="_self">Policy</a></td><td class="desc">Internal structure of iterator - made public to enable introspection </td></tr>
<tr id="row_0_3_4_16_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span id="arr_0_3_4_16_" class="arrow" onclick="toggleFolder('0_3_4_16_')">&#9658;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1gemm_1_1warp_1_1MmaTensorOpAccumulatorTileIterator_3_01Shape___00_01Element___006c39f57875e0aa9d0ad82c8043ed8b98.html" target="_self">MmaTensorOpAccumulatorTileIterator&lt; Shape_, Element_, cutlass::layout::RowMajor, InstructionShape_, OpDelta_ &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_3_4_16_0_" style="display:none;"><td class="entry"><span style="width:80px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1gemm_1_1warp_1_1MmaTensorOpAccumulatorTileIterator_3_01Shape___00_01Element___093b5d2838ac5a742704ef62b5c8688f0.html" target="_self">Policy</a></td><td class="desc">Internal structure of iterator - made public to enable introspection </td></tr>
<tr id="row_0_3_4_17_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1gemm_1_1warp_1_1MmaTensorOpMultiplicandTileIterator.html" target="_self">MmaTensorOpMultiplicandTileIterator</a></td><td class="desc"></td></tr>
<tr id="row_0_3_4_18_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1gemm_1_1warp_1_1MmaTensorOpMultiplicandTileIterator_3_01Shape___00_01Operand___0b84f53cd44b339eccc12067c9f86e11c.html" target="_self">MmaTensorOpMultiplicandTileIterator&lt; Shape_, Operand_, Element_, cutlass::layout::ColumnMajorTensorOpMultiplicandCongruous&lt; sizeof_bits&lt; Element_ &gt;::value, int(128/sizeof(Element_))&gt;, InstructionShape_, OpDelta_, 32, PartitionsK_ &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_3_4_19_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1gemm_1_1warp_1_1MmaTensorOpMultiplicandTileIterator_3_01Shape___00_01Operand___0e52ad425e1ee3e68544873f66733237b.html" target="_self">MmaTensorOpMultiplicandTileIterator&lt; Shape_, Operand_, Element_, cutlass::layout::ColumnMajorTensorOpMultiplicandCrosswise&lt; sizeof_bits&lt; Element_ &gt;::value, Crosswise &gt;, InstructionShape_, OpDelta_, 32, PartitionsK_ &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_3_4_20_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1gemm_1_1warp_1_1MmaTensorOpMultiplicandTileIterator_3_01Shape___00_01Operand___039819fb3ccd43786d556c2c9669508ef.html" target="_self">MmaTensorOpMultiplicandTileIterator&lt; Shape_, Operand_, Element_, cutlass::layout::RowMajorTensorOpMultiplicandCongruous&lt; sizeof_bits&lt; Element_ &gt;::value, int(128/sizeof(Element_))&gt;, InstructionShape_, OpDelta_, 32, PartitionsK_ &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_3_4_21_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1gemm_1_1warp_1_1MmaTensorOpMultiplicandTileIterator_3_01Shape___00_01Operand___0352e0dcab42bc8360606874e00173556.html" target="_self">MmaTensorOpMultiplicandTileIterator&lt; Shape_, Operand_, Element_, cutlass::layout::RowMajorTensorOpMultiplicandCrosswise&lt; sizeof_bits&lt; Element_ &gt;::value, Crosswise &gt;, InstructionShape_, OpDelta_, 32, PartitionsK_ &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_3_4_22_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span id="arr_0_3_4_22_" class="arrow" onclick="toggleFolder('0_3_4_22_')">&#9658;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1gemm_1_1warp_1_1MmaTensorOpMultiplicandTileIterator_3_01Shape___00_01Operand___0ed7daaeba1c095e77f68533d4d2c475c.html" target="_self">MmaTensorOpMultiplicandTileIterator&lt; Shape_, Operand_, Element_, cutlass::layout::TensorOpMultiplicandCongruous&lt; sizeof_bits&lt; Element_ &gt;::value, 64 &gt;, InstructionShape_, OpDelta_, 32, PartitionsK_ &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_3_4_22_0_" style="display:none;"><td class="entry"><span style="width:80px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1gemm_1_1warp_1_1MmaTensorOpMultiplicandTileIterator_3_01Shape___00_01Operand___07638f8b7761f6e2e2e6918e2c05e739.html" target="_self">Policy</a></td><td class="desc">Internal structure of iterator - made public to enable introspection </td></tr>
<tr id="row_0_3_4_23_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span id="arr_0_3_4_23_" class="arrow" onclick="toggleFolder('0_3_4_23_')">&#9658;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1gemm_1_1warp_1_1MmaTensorOpMultiplicandTileIterator_3_01Shape___00_01Operand___0c7d419c589d601ce4eb603be566fea21.html" target="_self">MmaTensorOpMultiplicandTileIterator&lt; Shape_, Operand_, Element_, cutlass::layout::TensorOpMultiplicandCrosswise&lt; sizeof_bits&lt; Element_ &gt;::value, Crosswise &gt;, InstructionShape_, OpDelta_, 32, PartitionsK_ &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_3_4_23_0_" style="display:none;"><td class="entry"><span style="width:80px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1gemm_1_1warp_1_1MmaTensorOpMultiplicandTileIterator_3_01Shape___00_01Operand___0784c74bd670999ec23ad8ef9dc55777.html" target="_self">Policy</a></td><td class="desc">Internal structure of iterator - made public to enable introspection </td></tr>
<tr id="row_0_3_4_24_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1gemm_1_1warp_1_1MmaTensorOpPolicy.html" target="_self">MmaTensorOpPolicy</a></td><td class="desc">Policy </td></tr>
<tr id="row_0_3_4_25_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1gemm_1_1warp_1_1MmaVoltaTensorOp.html" target="_self">MmaVoltaTensorOp</a></td><td class="desc">Structure to compute the matrix product targeting CUDA cores and SIMT math instructions </td></tr>
<tr id="row_0_3_4_26_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span id="arr_0_3_4_26_" class="arrow" onclick="toggleFolder('0_3_4_26_')">&#9658;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1gemm_1_1warp_1_1MmaVoltaTensorOpAccumulatorTileIterator.html" target="_self">MmaVoltaTensorOpAccumulatorTileIterator</a></td><td class="desc"></td></tr>
<tr id="row_0_3_4_26_0_" style="display:none;"><td class="entry"><span style="width:80px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1gemm_1_1warp_1_1MmaVoltaTensorOpAccumulatorTileIterator_1_1Policy.html" target="_self">Policy</a></td><td class="desc">Internal structure of iterator - made public to enable introspection </td></tr>
<tr id="row_0_3_4_27_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1gemm_1_1warp_1_1MmaVoltaTensorOpMultiplicandTileIterator.html" target="_self">MmaVoltaTensorOpMultiplicandTileIterator</a></td><td class="desc"></td></tr>
<tr id="row_0_3_4_28_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1gemm_1_1warp_1_1MmaVoltaTensorOpMultiplicandTileIterator_3_01Shape___00_01Operan0d3248553e52cd61ed8a2b3b12a20343.html" target="_self">MmaVoltaTensorOpMultiplicandTileIterator&lt; Shape_, Operand::kA, Element_, cutlass::layout::ColumnMajorVoltaTensorOpMultiplicandCongruous&lt; sizeof_bits&lt; Element_ &gt;::value &gt;, InstructionShape_, OpDelta_, 32 &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_3_4_29_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span id="arr_0_3_4_29_" class="arrow" onclick="toggleFolder('0_3_4_29_')">&#9658;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1gemm_1_1warp_1_1MmaVoltaTensorOpMultiplicandTileIterator_3_01Shape___00_01Operan34be8e21a40af3ebd2dc3dff460dca72.html" target="_self">MmaVoltaTensorOpMultiplicandTileIterator&lt; Shape_, Operand::kA, Element_, cutlass::layout::VoltaTensorOpMultiplicandCongruous&lt; sizeof_bits&lt; Element_ &gt;::value &gt;, InstructionShape_, OpDelta_, 32 &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_3_4_29_0_" style="display:none;"><td class="entry"><span style="width:80px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1gemm_1_1warp_1_1MmaVoltaTensorOpMultiplicandTileIterator_3_01Shape___00_01Opera33cdf53848564e894d4407637dc86caf.html" target="_self">Policy</a></td><td class="desc">Internal structure of iterator - made public to enable introspection </td></tr>
<tr id="row_0_3_4_30_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1gemm_1_1warp_1_1MmaVoltaTensorOpMultiplicandTileIterator_3_01Shape___00_01Operand734577b7e54a074d143aba59828c2f2.html" target="_self">MmaVoltaTensorOpMultiplicandTileIterator&lt; Shape_, Operand::kB, Element_, cutlass::layout::RowMajorVoltaTensorOpMultiplicandBCongruous&lt; sizeof_bits&lt; Element_ &gt;::value &gt;, InstructionShape_, OpDelta_, 32 &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_3_4_31_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span id="arr_0_3_4_31_" class="arrow" onclick="toggleFolder('0_3_4_31_')">&#9658;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1gemm_1_1warp_1_1MmaVoltaTensorOpMultiplicandTileIterator_3_01Shape___00_01Operan16c56cdc2dda5eeb996af8ec0242d501.html" target="_self">MmaVoltaTensorOpMultiplicandTileIterator&lt; Shape_, Operand::kB, Element_, cutlass::layout::VoltaTensorOpMultiplicandBCongruous&lt; sizeof_bits&lt; Element_ &gt;::value &gt;, InstructionShape_, OpDelta_, 32 &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_3_4_31_0_" style="display:none;"><td class="entry"><span style="width:80px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1gemm_1_1warp_1_1MmaVoltaTensorOpMultiplicandTileIterator_3_01Shape___00_01Opera6fa6d2d3725bb3ec613d5c527ea3ffe7.html" target="_self">Policy</a></td><td class="desc">Internal structure of iterator - made public to enable introspection </td></tr>
<tr id="row_0_3_4_32_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1gemm_1_1warp_1_1MmaVoltaTensorOpMultiplicandTileIterator_3_01Shape___00_01Operan5a221944f4a0e16ccab77ba684856942.html" target="_self">MmaVoltaTensorOpMultiplicandTileIterator&lt; Shape_, Operand_, Element_, cutlass::layout::ColumnMajorVoltaTensorOpMultiplicandCrosswise&lt; sizeof_bits&lt; Element_ &gt;::value, KBlock &gt;, InstructionShape_, OpDelta_, 32 &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_3_4_33_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1gemm_1_1warp_1_1MmaVoltaTensorOpMultiplicandTileIterator_3_01Shape___00_01Operandcc9821c435540895138bc9af495f321.html" target="_self">MmaVoltaTensorOpMultiplicandTileIterator&lt; Shape_, Operand_, Element_, cutlass::layout::RowMajorVoltaTensorOpMultiplicandCrosswise&lt; sizeof_bits&lt; Element_ &gt;::value, KBlock &gt;, InstructionShape_, OpDelta_, 32 &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_3_4_34_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span id="arr_0_3_4_34_" class="arrow" onclick="toggleFolder('0_3_4_34_')">&#9658;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1gemm_1_1warp_1_1MmaVoltaTensorOpMultiplicandTileIterator_3_01Shape___00_01Operana2f40b28f0d2286b84d86f7238d67b52.html" target="_self">MmaVoltaTensorOpMultiplicandTileIterator&lt; Shape_, Operand_, Element_, cutlass::layout::VoltaTensorOpMultiplicandCrosswise&lt; sizeof_bits&lt; Element_ &gt;::value, KBlock &gt;, InstructionShape_, OpDelta_, 32 &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_3_4_34_0_" style="display:none;"><td class="entry"><span style="width:80px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1gemm_1_1warp_1_1MmaVoltaTensorOpMultiplicandTileIterator_3_01Shape___00_01Operafa294175b280756dd8388f9ffe7b72c4.html" target="_self">Policy</a></td><td class="desc">Internal structure of iterator - made public to enable introspection </td></tr>
<tr id="row_0_3_4_35_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1gemm_1_1warp_1_1WarpSize.html" target="_self">WarpSize</a></td><td class="desc">Query the number of threads per warp </td></tr>
<tr id="row_0_3_5_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1gemm_1_1BatchedGemmCoord.html" target="_self">BatchedGemmCoord</a></td><td class="desc"></td></tr>
<tr id="row_0_3_6_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1gemm_1_1GemmCoord.html" target="_self">GemmCoord</a></td><td class="desc"></td></tr>
<tr id="row_0_3_7_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1gemm_1_1GemmShape.html" target="_self">GemmShape</a></td><td class="desc">Shape of a matrix multiply-add operation </td></tr>
<tr id="row_0_4_" style="display:none;"><td class="entry"><span style="width:16px;display:inline-block;">&#160;</span><span id="arr_0_4_" class="arrow" onclick="toggleFolder('0_4_')">&#9658;</span><span class="icona"><span class="icon">N</span></span><a class="el" href="namespacecutlass_1_1layout.html" target="_self">layout</a></td><td class="desc"></td></tr>
<tr id="row_0_4_0_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1layout_1_1ColumnMajor.html" target="_self">ColumnMajor</a></td><td class="desc">Mapping function for column-major matrices </td></tr>
<tr id="row_0_4_1_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1layout_1_1ColumnMajorBlockLinear.html" target="_self">ColumnMajorBlockLinear</a></td><td class="desc"></td></tr>
<tr id="row_0_4_2_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1layout_1_1ColumnMajorInterleaved.html" target="_self">ColumnMajorInterleaved</a></td><td class="desc"></td></tr>
<tr id="row_0_4_3_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1layout_1_1ColumnMajorTensorOpMultiplicandCongruous.html" target="_self">ColumnMajorTensorOpMultiplicandCongruous</a></td><td class="desc"></td></tr>
<tr id="row_0_4_4_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1layout_1_1ColumnMajorTensorOpMultiplicandCrosswise.html" target="_self">ColumnMajorTensorOpMultiplicandCrosswise</a></td><td class="desc"></td></tr>
<tr id="row_0_4_5_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1layout_1_1ColumnMajorVoltaTensorOpMultiplicandBCongruous.html" target="_self">ColumnMajorVoltaTensorOpMultiplicandBCongruous</a></td><td class="desc">Template mapping a column-major view of pitch-linear memory to <a class="el" href="structcutlass_1_1layout_1_1VoltaTensorOpMultiplicandCongruous.html" title="Template based on element size (in bits) - defined in terms of pitch-linear memory. ">VoltaTensorOpMultiplicandCongruous</a> </td></tr>
<tr id="row_0_4_6_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1layout_1_1ColumnMajorVoltaTensorOpMultiplicandCongruous.html" target="_self">ColumnMajorVoltaTensorOpMultiplicandCongruous</a></td><td class="desc">Template mapping a column-major view of pitch-linear memory to <a class="el" href="structcutlass_1_1layout_1_1VoltaTensorOpMultiplicandCongruous.html" title="Template based on element size (in bits) - defined in terms of pitch-linear memory. ">VoltaTensorOpMultiplicandCongruous</a> </td></tr>
<tr id="row_0_4_7_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1layout_1_1ColumnMajorVoltaTensorOpMultiplicandCrosswise.html" target="_self">ColumnMajorVoltaTensorOpMultiplicandCrosswise</a></td><td class="desc"></td></tr>
<tr id="row_0_4_8_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1layout_1_1ContiguousMatrix.html" target="_self">ContiguousMatrix</a></td><td class="desc"></td></tr>
<tr id="row_0_4_9_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1layout_1_1GeneralMatrix.html" target="_self">GeneralMatrix</a></td><td class="desc"></td></tr>
<tr id="row_0_4_10_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1layout_1_1LayoutTranspose.html" target="_self">LayoutTranspose</a></td><td class="desc">Defines transposes of matrix layouts </td></tr>
<tr id="row_0_4_11_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1layout_1_1LayoutTranspose_3_01layout_1_1ColumnMajor_01_4.html" target="_self">LayoutTranspose&lt; layout::ColumnMajor &gt;</a></td><td class="desc">Transpose of column-major is row-major </td></tr>
<tr id="row_0_4_12_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1layout_1_1LayoutTranspose_3_01layout_1_1RowMajor_01_4.html" target="_self">LayoutTranspose&lt; layout::RowMajor &gt;</a></td><td class="desc">Transpose of row-major is column-major </td></tr>
<tr id="row_0_4_13_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1layout_1_1PackedVectorLayout.html" target="_self">PackedVectorLayout</a></td><td class="desc">Tensor layout for densely packed vectors </td></tr>
<tr id="row_0_4_14_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1layout_1_1PitchLinear.html" target="_self">PitchLinear</a></td><td class="desc">Mapping function for pitch-linear memory </td></tr>
<tr id="row_0_4_15_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1layout_1_1PitchLinearCoord.html" target="_self">PitchLinearCoord</a></td><td class="desc">Coordinate in pitch-linear space </td></tr>
<tr id="row_0_4_16_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1layout_1_1PitchLinearShape.html" target="_self">PitchLinearShape</a></td><td class="desc">Template defining a shape used by pitch-linear operators </td></tr>
<tr id="row_0_4_17_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1layout_1_1RowMajor.html" target="_self">RowMajor</a></td><td class="desc">Mapping function for row-major matrices </td></tr>
<tr id="row_0_4_18_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1layout_1_1RowMajorBlockLinear.html" target="_self">RowMajorBlockLinear</a></td><td class="desc"></td></tr>
<tr id="row_0_4_19_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1layout_1_1RowMajorInterleaved.html" target="_self">RowMajorInterleaved</a></td><td class="desc"></td></tr>
<tr id="row_0_4_20_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1layout_1_1RowMajorTensorOpMultiplicandCongruous.html" target="_self">RowMajorTensorOpMultiplicandCongruous</a></td><td class="desc"></td></tr>
<tr id="row_0_4_21_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1layout_1_1RowMajorTensorOpMultiplicandCrosswise.html" target="_self">RowMajorTensorOpMultiplicandCrosswise</a></td><td class="desc"></td></tr>
<tr id="row_0_4_22_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1layout_1_1RowMajorVoltaTensorOpMultiplicandBCongruous.html" target="_self">RowMajorVoltaTensorOpMultiplicandBCongruous</a></td><td class="desc">Template mapping a row-major view of pitch-linear memory to <a class="el" href="structcutlass_1_1layout_1_1VoltaTensorOpMultiplicandCongruous.html" title="Template based on element size (in bits) - defined in terms of pitch-linear memory. ">VoltaTensorOpMultiplicandCongruous</a> </td></tr>
<tr id="row_0_4_23_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1layout_1_1RowMajorVoltaTensorOpMultiplicandCongruous.html" target="_self">RowMajorVoltaTensorOpMultiplicandCongruous</a></td><td class="desc">Template mapping a row-major view of pitch-linear memory to <a class="el" href="structcutlass_1_1layout_1_1VoltaTensorOpMultiplicandCongruous.html" title="Template based on element size (in bits) - defined in terms of pitch-linear memory. ">VoltaTensorOpMultiplicandCongruous</a> </td></tr>
<tr id="row_0_4_24_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1layout_1_1RowMajorVoltaTensorOpMultiplicandCrosswise.html" target="_self">RowMajorVoltaTensorOpMultiplicandCrosswise</a></td><td class="desc"></td></tr>
<tr id="row_0_4_25_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1layout_1_1TensorCxRSKx.html" target="_self">TensorCxRSKx</a></td><td class="desc">Mapping function for 4-D CxRSKx tensors </td></tr>
<tr id="row_0_4_26_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1layout_1_1TensorNCHW.html" target="_self">TensorNCHW</a></td><td class="desc">Mapping function for 4-D NCHW tensors </td></tr>
<tr id="row_0_4_27_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1layout_1_1TensorNCxHWx.html" target="_self">TensorNCxHWx</a></td><td class="desc">Mapping function for 4-D NC/xHWx tensors </td></tr>
<tr id="row_0_4_28_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1layout_1_1TensorNHWC.html" target="_self">TensorNHWC</a></td><td class="desc">Mapping function for 4-D NHWC tensors </td></tr>
<tr id="row_0_4_29_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1layout_1_1TensorOpMultiplicand.html" target="_self">TensorOpMultiplicand</a></td><td class="desc"></td></tr>
<tr id="row_0_4_30_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1layout_1_1TensorOpMultiplicandColumnMajorInterleaved.html" target="_self">TensorOpMultiplicandColumnMajorInterleaved</a></td><td class="desc">Template based on element size (in bits) - defined in terms of pitch-linear memory </td></tr>
<tr id="row_0_4_31_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1layout_1_1TensorOpMultiplicandCongruous.html" target="_self">TensorOpMultiplicandCongruous</a></td><td class="desc"></td></tr>
<tr id="row_0_4_32_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1layout_1_1TensorOpMultiplicandCongruous_3_0132_00_01Crosswise_01_4.html" target="_self">TensorOpMultiplicandCongruous&lt; 32, Crosswise &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_4_33_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1layout_1_1TensorOpMultiplicandCrosswise.html" target="_self">TensorOpMultiplicandCrosswise</a></td><td class="desc"></td></tr>
<tr id="row_0_4_34_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1layout_1_1TensorOpMultiplicandRowMajorInterleaved.html" target="_self">TensorOpMultiplicandRowMajorInterleaved</a></td><td class="desc">Template based on element size (in bits) - defined in terms of pitch-linear memory </td></tr>
<tr id="row_0_4_35_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1layout_1_1VoltaTensorOpMultiplicandBCongruous.html" target="_self">VoltaTensorOpMultiplicandBCongruous</a></td><td class="desc">Template based on element size (in bits) - defined in terms of pitch-linear memory </td></tr>
<tr id="row_0_4_36_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1layout_1_1VoltaTensorOpMultiplicandCongruous.html" target="_self">VoltaTensorOpMultiplicandCongruous</a></td><td class="desc">Template based on element size (in bits) - defined in terms of pitch-linear memory </td></tr>
<tr id="row_0_4_37_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1layout_1_1VoltaTensorOpMultiplicandCrosswise.html" target="_self">VoltaTensorOpMultiplicandCrosswise</a></td><td class="desc"></td></tr>
<tr id="row_0_5_" style="display:none;"><td class="entry"><span style="width:16px;display:inline-block;">&#160;</span><span id="arr_0_5_" class="arrow" onclick="toggleFolder('0_5_')">&#9658;</span><span class="icona"><span class="icon">N</span></span><a class="el" href="namespacecutlass_1_1library.html" target="_self">library</a></td><td class="desc"></td></tr>
<tr id="row_0_5_0_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1library_1_1GemmArguments.html" target="_self">GemmArguments</a></td><td class="desc">Arguments for GEMM </td></tr>
<tr id="row_0_5_1_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1library_1_1GemmArrayArguments.html" target="_self">GemmArrayArguments</a></td><td class="desc">Arguments for GEMM - used by all the GEMM operations </td></tr>
<tr id="row_0_5_2_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1library_1_1GemmArrayConfiguration.html" target="_self">GemmArrayConfiguration</a></td><td class="desc">Configuration for batched GEMM in which multiple matrix products are computed </td></tr>
<tr id="row_0_5_3_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1library_1_1GemmBatchedConfiguration.html" target="_self">GemmBatchedConfiguration</a></td><td class="desc">Configuration for batched GEMM in which multiple matrix products are computed </td></tr>
<tr id="row_0_5_4_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1library_1_1GemmConfiguration.html" target="_self">GemmConfiguration</a></td><td class="desc">Configuration for basic GEMM operations </td></tr>
<tr id="row_0_5_5_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1library_1_1GemmDescription.html" target="_self">GemmDescription</a></td><td class="desc">Description of all GEMM computations </td></tr>
<tr id="row_0_5_6_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1library_1_1GemmPlanarComplexBatchedConfiguration.html" target="_self">GemmPlanarComplexBatchedConfiguration</a></td><td class="desc">Batched complex valued GEMM in which real and imaginary parts are separated by a stride </td></tr>
<tr id="row_0_5_7_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1library_1_1GemmPlanarComplexConfiguration.html" target="_self">GemmPlanarComplexConfiguration</a></td><td class="desc">Complex valued GEMM in which real and imaginary parts are separated by a stride </td></tr>
<tr id="row_0_5_8_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1library_1_1Manifest.html" target="_self">Manifest</a></td><td class="desc"><a class="el" href="classcutlass_1_1library_1_1Manifest.html" title="Manifest of CUTLASS Library. ">Manifest</a> of CUTLASS Library </td></tr>
<tr id="row_0_5_9_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1library_1_1MathInstructionDescription.html" target="_self">MathInstructionDescription</a></td><td class="desc"></td></tr>
<tr id="row_0_5_10_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1library_1_1Operation.html" target="_self">Operation</a></td><td class="desc">Base class for all device-wide operations </td></tr>
<tr id="row_0_5_11_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1library_1_1OperationDescription.html" target="_self">OperationDescription</a></td><td class="desc">High-level description of an operation </td></tr>
<tr id="row_0_5_12_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1library_1_1TensorDescription.html" target="_self">TensorDescription</a></td><td class="desc">Structure describing the properties of a tensor </td></tr>
<tr id="row_0_5_13_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1library_1_1TileDescription.html" target="_self">TileDescription</a></td><td class="desc">Structure describing the tiled structure of a GEMM-like computation </td></tr>
<tr id="row_0_6_" style="display:none;"><td class="entry"><span style="width:16px;display:inline-block;">&#160;</span><span id="arr_0_6_" class="arrow" onclick="toggleFolder('0_6_')">&#9658;</span><span class="icona"><span class="icon">N</span></span><a class="el" href="namespacecutlass_1_1platform.html" target="_self">platform</a></td><td class="desc"></td></tr>
<tr id="row_0_6_0_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1platform_1_1aligned__chunk.html" target="_self">aligned_chunk</a></td><td class="desc"></td></tr>
<tr id="row_0_6_1_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1platform_1_1aligned__storage.html" target="_self">aligned_storage</a></td><td class="desc">Std::aligned_storage </td></tr>
<tr id="row_0_6_2_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span id="arr_0_6_2_" class="arrow" onclick="toggleFolder('0_6_2_')">&#9658;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1platform_1_1alignment__of.html" target="_self">alignment_of</a></td><td class="desc">Std::alignment_of </td></tr>
<tr id="row_0_6_2_0_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1platform_1_1alignment__of_1_1pad.html" target="_self">pad</a></td><td class="desc"></td></tr>
<tr id="row_0_6_3_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1platform_1_1alignment__of_3_01const_01value__t_01_4.html" target="_self">alignment_of&lt; const value_t &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_6_4_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1platform_1_1alignment__of_3_01const_01volatile_01value__t_01_4.html" target="_self">alignment_of&lt; const volatile value_t &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_6_5_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1platform_1_1alignment__of_3_01double2_01_4.html" target="_self">alignment_of&lt; double2 &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_6_6_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1platform_1_1alignment__of_3_01double4_01_4.html" target="_self">alignment_of&lt; double4 &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_6_7_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1platform_1_1alignment__of_3_01float4_01_4.html" target="_self">alignment_of&lt; float4 &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_6_8_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1platform_1_1alignment__of_3_01int4_01_4.html" target="_self">alignment_of&lt; int4 &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_6_9_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1platform_1_1alignment__of_3_01long4_01_4.html" target="_self">alignment_of&lt; long4 &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_6_10_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1platform_1_1alignment__of_3_01longlong2_01_4.html" target="_self">alignment_of&lt; longlong2 &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_6_11_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1platform_1_1alignment__of_3_01longlong4_01_4.html" target="_self">alignment_of&lt; longlong4 &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_6_12_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1platform_1_1alignment__of_3_01uint4_01_4.html" target="_self">alignment_of&lt; uint4 &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_6_13_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1platform_1_1alignment__of_3_01ulong4_01_4.html" target="_self">alignment_of&lt; ulong4 &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_6_14_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1platform_1_1alignment__of_3_01ulonglong2_01_4.html" target="_self">alignment_of&lt; ulonglong2 &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_6_15_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1platform_1_1alignment__of_3_01ulonglong4_01_4.html" target="_self">alignment_of&lt; ulonglong4 &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_6_16_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1platform_1_1alignment__of_3_01volatile_01value__t_01_4.html" target="_self">alignment_of&lt; volatile value_t &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_6_17_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1platform_1_1bool__constant.html" target="_self">bool_constant</a></td><td class="desc">Std::bool_constant </td></tr>
<tr id="row_0_6_18_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1platform_1_1conditional.html" target="_self">conditional</a></td><td class="desc">Std::conditional (true specialization) </td></tr>
<tr id="row_0_6_19_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1platform_1_1conditional_3_01false_00_01T_00_01F_01_4.html" target="_self">conditional&lt; false, T, F &gt;</a></td><td class="desc">Std::conditional (false specialization) </td></tr>
<tr id="row_0_6_20_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1platform_1_1default__delete.html" target="_self">default_delete</a></td><td class="desc">Default deleter </td></tr>
<tr id="row_0_6_21_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1platform_1_1default__delete_3_01T[]_4.html" target="_self">default_delete&lt; T[]&gt;</a></td><td class="desc">Partial specialization for deleting array types </td></tr>
<tr id="row_0_6_22_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1platform_1_1enable__if.html" target="_self">enable_if</a></td><td class="desc">Std::enable_if (true specialization) </td></tr>
<tr id="row_0_6_23_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1platform_1_1enable__if_3_01false_00_01T_01_4.html" target="_self">enable_if&lt; false, T &gt;</a></td><td class="desc">Std::enable_if (false specialization) </td></tr>
<tr id="row_0_6_24_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1platform_1_1integral__constant.html" target="_self">integral_constant</a></td><td class="desc">Std::integral_constant </td></tr>
<tr id="row_0_6_25_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1platform_1_1is__arithmetic.html" target="_self">is_arithmetic</a></td><td class="desc">Std::is_arithmetic </td></tr>
<tr id="row_0_6_26_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1platform_1_1is__base__of.html" target="_self">is_base_of</a></td><td class="desc">Std::is_base_of </td></tr>
<tr id="row_0_6_27_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span id="arr_0_6_27_" class="arrow" onclick="toggleFolder('0_6_27_')">&#9658;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1platform_1_1is__base__of__helper.html" target="_self">is_base_of_helper</a></td><td class="desc">Helper for std::is_base_of </td></tr>
<tr id="row_0_6_27_0_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1platform_1_1is__base__of__helper_1_1dummy.html" target="_self">dummy</a></td><td class="desc"></td></tr>
<tr id="row_0_6_28_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1platform_1_1is__floating__point.html" target="_self">is_floating_point</a></td><td class="desc">Std::is_floating_point </td></tr>
<tr id="row_0_6_29_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1platform_1_1is__fundamental.html" target="_self">is_fundamental</a></td><td class="desc">Std::is_fundamental </td></tr>
<tr id="row_0_6_30_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1platform_1_1is__integral.html" target="_self">is_integral</a></td><td class="desc">Std::is_integral </td></tr>
<tr id="row_0_6_31_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1platform_1_1is__integral_3_01char_01_4.html" target="_self">is_integral&lt; char &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_6_32_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1platform_1_1is__integral_3_01const_01T_01_4.html" target="_self">is_integral&lt; const T &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_6_33_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1platform_1_1is__integral_3_01const_01volatile_01T_01_4.html" target="_self">is_integral&lt; const volatile T &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_6_34_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1platform_1_1is__integral_3_01int_01_4.html" target="_self">is_integral&lt; int &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_6_35_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1platform_1_1is__integral_3_01long_01_4.html" target="_self">is_integral&lt; long &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_6_36_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1platform_1_1is__integral_3_01long_01long_01_4.html" target="_self">is_integral&lt; long long &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_6_37_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1platform_1_1is__integral_3_01short_01_4.html" target="_self">is_integral&lt; short &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_6_38_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1platform_1_1is__integral_3_01signed_01char_01_4.html" target="_self">is_integral&lt; signed char &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_6_39_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1platform_1_1is__integral_3_01unsigned_01char_01_4.html" target="_self">is_integral&lt; unsigned char &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_6_40_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1platform_1_1is__integral_3_01unsigned_01int_01_4.html" target="_self">is_integral&lt; unsigned int &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_6_41_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1platform_1_1is__integral_3_01unsigned_01long_01_4.html" target="_self">is_integral&lt; unsigned long &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_6_42_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1platform_1_1is__integral_3_01unsigned_01long_01long_01_4.html" target="_self">is_integral&lt; unsigned long long &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_6_43_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1platform_1_1is__integral_3_01unsigned_01short_01_4.html" target="_self">is_integral&lt; unsigned short &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_6_44_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1platform_1_1is__integral_3_01volatile_01T_01_4.html" target="_self">is_integral&lt; volatile T &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_6_45_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1platform_1_1is__pointer.html" target="_self">is_pointer</a></td><td class="desc">Std::is_pointer </td></tr>
<tr id="row_0_6_46_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1platform_1_1is__pointer__helper.html" target="_self">is_pointer_helper</a></td><td class="desc">Helper for std::is_pointer (false specialization) </td></tr>
<tr id="row_0_6_47_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1platform_1_1is__pointer__helper_3_01T_01_5_01_4.html" target="_self">is_pointer_helper&lt; T * &gt;</a></td><td class="desc">Helper for std::is_pointer (true specialization) </td></tr>
<tr id="row_0_6_48_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1platform_1_1is__same.html" target="_self">is_same</a></td><td class="desc">Std::is_same (false specialization) </td></tr>
<tr id="row_0_6_49_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1platform_1_1is__same_3_01A_00_01A_01_4.html" target="_self">is_same&lt; A, A &gt;</a></td><td class="desc">Std::is_same (true specialization) </td></tr>
<tr id="row_0_6_50_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1platform_1_1is__trivially__copyable.html" target="_self">is_trivially_copyable</a></td><td class="desc"></td></tr>
<tr id="row_0_6_51_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1platform_1_1is__void.html" target="_self">is_void</a></td><td class="desc">Std::is_void </td></tr>
<tr id="row_0_6_52_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1platform_1_1is__volatile.html" target="_self">is_volatile</a></td><td class="desc">Std::is_volatile </td></tr>
<tr id="row_0_6_53_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1platform_1_1is__volatile_3_01volatile_01T_01_4.html" target="_self">is_volatile&lt; volatile T &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_6_54_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1platform_1_1nullptr__t.html" target="_self">nullptr_t</a></td><td class="desc">Std::nullptr_t </td></tr>
<tr id="row_0_6_55_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1platform_1_1remove__const.html" target="_self">remove_const</a></td><td class="desc">Std::remove_const (non-const specialization) </td></tr>
<tr id="row_0_6_56_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1platform_1_1remove__const_3_01const_01T_01_4.html" target="_self">remove_const&lt; const T &gt;</a></td><td class="desc">Std::remove_const (const specialization) </td></tr>
<tr id="row_0_6_57_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1platform_1_1remove__cv.html" target="_self">remove_cv</a></td><td class="desc">Std::remove_cv </td></tr>
<tr id="row_0_6_58_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1platform_1_1remove__volatile.html" target="_self">remove_volatile</a></td><td class="desc">Std::remove_volatile (non-volatile specialization) </td></tr>
<tr id="row_0_6_59_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1platform_1_1remove__volatile_3_01volatile_01T_01_4.html" target="_self">remove_volatile&lt; volatile T &gt;</a></td><td class="desc">Std::remove_volatile (volatile specialization) </td></tr>
<tr id="row_0_6_60_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1platform_1_1unique__ptr.html" target="_self">unique_ptr</a></td><td class="desc">Std::unique_ptr </td></tr>
<tr id="row_0_7_" style="display:none;"><td class="entry"><span style="width:16px;display:inline-block;">&#160;</span><span id="arr_0_7_" class="arrow" onclick="toggleFolder('0_7_')">&#9658;</span><span class="icona"><span class="icon">N</span></span><a class="el" href="namespacecutlass_1_1reduction.html" target="_self">reduction</a></td><td class="desc"></td></tr>
<tr id="row_0_7_0_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span id="arr_0_7_0_" class="arrow" onclick="toggleFolder('0_7_0_')">&#9658;</span><span class="icona"><span class="icon">N</span></span><a class="el" href="namespacecutlass_1_1reduction_1_1kernel.html" target="_self">kernel</a></td><td class="desc"></td></tr>
<tr id="row_0_7_0_0_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span id="arr_0_7_0_0_" class="arrow" onclick="toggleFolder('0_7_0_0_')">&#9658;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1reduction_1_1kernel_1_1ReduceSplitK.html" target="_self">ReduceSplitK</a></td><td class="desc"></td></tr>
<tr id="row_0_7_0_0_0_" style="display:none;"><td class="entry"><span style="width:80px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1reduction_1_1kernel_1_1ReduceSplitK_1_1Params.html" target="_self">Params</a></td><td class="desc"><a class="el" href="structcutlass_1_1reduction_1_1kernel_1_1ReduceSplitK_1_1Params.html" title="Params structure. ">Params</a> structure </td></tr>
<tr id="row_0_7_0_0_1_" style="display:none;"><td class="entry"><span style="width:80px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1reduction_1_1kernel_1_1ReduceSplitK_1_1SharedStorage.html" target="_self">SharedStorage</a></td><td class="desc"></td></tr>
<tr id="row_0_7_1_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span id="arr_0_7_1_" class="arrow" onclick="toggleFolder('0_7_1_')">&#9658;</span><span class="icona"><span class="icon">N</span></span><a class="el" href="namespacecutlass_1_1reduction_1_1thread.html" target="_self">thread</a></td><td class="desc"></td></tr>
<tr id="row_0_7_1_0_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1reduction_1_1thread_1_1Reduce.html" target="_self">Reduce</a></td><td class="desc">Structure to compute the thread level reduction </td></tr>
<tr id="row_0_7_1_1_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1reduction_1_1thread_1_1Reduce_3_01plus_3_01half__t_01_4_00_01AlignedArray_3_01half__t_00_01N_01_4_01_4.html" target="_self">Reduce&lt; plus&lt; half_t &gt;, AlignedArray&lt; half_t, N &gt; &gt;</a></td><td class="desc">Partial specializations of <a class="el" href="structcutlass_1_1reduction_1_1thread_1_1Reduce.html" title="Structure to compute the thread level reduction. ">Reduce</a> for AlignedArray&lt;half_t, N&gt; </td></tr>
<tr id="row_0_7_1_2_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1reduction_1_1thread_1_1Reduce_3_01plus_3_01half__t_01_4_00_01Array_3_01half__t_00_01N_01_4_01_4.html" target="_self">Reduce&lt; plus&lt; half_t &gt;, Array&lt; half_t, N &gt; &gt;</a></td><td class="desc">Partial specializations of <a class="el" href="structcutlass_1_1reduction_1_1thread_1_1Reduce.html" title="Structure to compute the thread level reduction. ">Reduce</a> for Array&lt;half_t, N&gt; </td></tr>
<tr id="row_0_7_1_3_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1reduction_1_1thread_1_1Reduce_3_01plus_3_01T_01_4_00_01Array_3_01T_00_01N_01_4_01_4.html" target="_self">Reduce&lt; plus&lt; T &gt;, Array&lt; T, N &gt; &gt;</a></td><td class="desc">Partial specialization of <a class="el" href="structcutlass_1_1reduction_1_1thread_1_1Reduce.html" title="Structure to compute the thread level reduction. ">Reduce</a> for Array&lt;T, N&gt; </td></tr>
<tr id="row_0_7_1_4_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1reduction_1_1thread_1_1Reduce_3_01plus_3_01T_01_4_00_01T_01_4.html" target="_self">Reduce&lt; plus&lt; T &gt;, T &gt;</a></td><td class="desc">Partial Specialization of <a class="el" href="structcutlass_1_1reduction_1_1thread_1_1Reduce.html" title="Structure to compute the thread level reduction. ">Reduce</a> for "plus" (a functional operator) </td></tr>
<tr id="row_0_7_1_5_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span id="arr_0_7_1_5_" class="arrow" onclick="toggleFolder('0_7_1_5_')">&#9658;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1reduction_1_1thread_1_1ReduceAdd.html" target="_self">ReduceAdd</a></td><td class="desc">Mixed-precision reduction </td></tr>
<tr id="row_0_7_1_5_0_" style="display:none;"><td class="entry"><span style="width:80px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1reduction_1_1thread_1_1ReduceAdd_1_1Params.html" target="_self">Params</a></td><td class="desc"></td></tr>
<tr id="row_0_7_2_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1reduction_1_1BatchedReduction.html" target="_self">BatchedReduction</a></td><td class="desc"></td></tr>
<tr id="row_0_7_3_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span id="arr_0_7_3_" class="arrow" onclick="toggleFolder('0_7_3_')">&#9658;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1reduction_1_1BatchedReductionTraits.html" target="_self">BatchedReductionTraits</a></td><td class="desc"></td></tr>
<tr id="row_0_7_3_0_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1reduction_1_1BatchedReductionTraits_1_1Params.html" target="_self">Params</a></td><td class="desc"></td></tr>
<tr id="row_0_7_4_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1reduction_1_1DefaultBlockSwizzle.html" target="_self">DefaultBlockSwizzle</a></td><td class="desc"></td></tr>
<tr id="row_0_8_" style="display:none;"><td class="entry"><span style="width:16px;display:inline-block;">&#160;</span><span id="arr_0_8_" class="arrow" onclick="toggleFolder('0_8_')">&#9658;</span><span class="icona"><span class="icon">N</span></span><a class="el" href="namespacecutlass_1_1reference.html" target="_self">reference</a></td><td class="desc"></td></tr>
<tr id="row_0_8_0_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span id="arr_0_8_0_" class="arrow" onclick="toggleFolder('0_8_0_')">&#9658;</span><span class="icona"><span class="icon">N</span></span><a class="el" href="namespacecutlass_1_1reference_1_1detail.html" target="_self">detail</a></td><td class="desc"></td></tr>
<tr id="row_0_8_0_0_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1reference_1_1detail_1_1Cast.html" target="_self">Cast</a></td><td class="desc"></td></tr>
<tr id="row_0_8_0_1_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1reference_1_1detail_1_1Cast_3_01float_00_01int8__t_01_4.html" target="_self">Cast&lt; float, int8_t &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_8_0_2_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1reference_1_1detail_1_1Cast_3_01float_00_01uint8__t_01_4.html" target="_self">Cast&lt; float, uint8_t &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_8_1_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span id="arr_0_8_1_" class="arrow" onclick="toggleFolder('0_8_1_')">&#9658;</span><span class="icona"><span class="icon">N</span></span><a class="el" href="namespacecutlass_1_1reference_1_1device.html" target="_self">device</a></td><td class="desc"></td></tr>
<tr id="row_0_8_1_0_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span id="arr_0_8_1_0_" class="arrow" onclick="toggleFolder('0_8_1_0_')">&#9658;</span><span class="icona"><span class="icon">N</span></span><a class="el" href="namespacecutlass_1_1reference_1_1device_1_1detail.html" target="_self">detail</a></td><td class="desc"></td></tr>
<tr id="row_0_8_1_0_0_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span id="arr_0_8_1_0_0_" class="arrow" onclick="toggleFolder('0_8_1_0_0_')">&#9658;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1reference_1_1device_1_1detail_1_1RandomGaussianFunc.html" target="_self">RandomGaussianFunc</a></td><td class="desc"></td></tr>
<tr id="row_0_8_1_0_0_0_" style="display:none;"><td class="entry"><span style="width:96px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1reference_1_1device_1_1detail_1_1RandomGaussianFunc_1_1Params.html" target="_self">Params</a></td><td class="desc">Parameters structure </td></tr>
<tr id="row_0_8_1_0_1_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span id="arr_0_8_1_0_1_" class="arrow" onclick="toggleFolder('0_8_1_0_1_')">&#9658;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1reference_1_1device_1_1detail_1_1RandomUniformFunc.html" target="_self">RandomUniformFunc</a></td><td class="desc">Computes a random Gaussian distribution </td></tr>
<tr id="row_0_8_1_0_1_0_" style="display:none;"><td class="entry"><span style="width:96px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1reference_1_1device_1_1detail_1_1RandomUniformFunc_1_1Params.html" target="_self">Params</a></td><td class="desc">Parameters structure </td></tr>
<tr id="row_0_8_1_0_2_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span id="arr_0_8_1_0_2_" class="arrow" onclick="toggleFolder('0_8_1_0_2_')">&#9658;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1reference_1_1device_1_1detail_1_1TensorCopyDiagonalInFunc.html" target="_self">TensorCopyDiagonalInFunc</a></td><td class="desc">Computes a random Gaussian distribution </td></tr>
<tr id="row_0_8_1_0_2_0_" style="display:none;"><td class="entry"><span style="width:96px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1reference_1_1device_1_1detail_1_1TensorCopyDiagonalInFunc_1_1Params.html" target="_self">Params</a></td><td class="desc">Parameters structure </td></tr>
<tr id="row_0_8_1_0_3_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span id="arr_0_8_1_0_3_" class="arrow" onclick="toggleFolder('0_8_1_0_3_')">&#9658;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1reference_1_1device_1_1detail_1_1TensorCopyDiagonalOutFunc.html" target="_self">TensorCopyDiagonalOutFunc</a></td><td class="desc">Computes a random Gaussian distribution </td></tr>
<tr id="row_0_8_1_0_3_0_" style="display:none;"><td class="entry"><span style="width:96px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1reference_1_1device_1_1detail_1_1TensorCopyDiagonalOutFunc_1_1Params.html" target="_self">Params</a></td><td class="desc">Parameters structure </td></tr>
<tr id="row_0_8_1_0_4_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span id="arr_0_8_1_0_4_" class="arrow" onclick="toggleFolder('0_8_1_0_4_')">&#9658;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1reference_1_1device_1_1detail_1_1TensorFillDiagonalFunc.html" target="_self">TensorFillDiagonalFunc</a></td><td class="desc">Computes a random Gaussian distribution </td></tr>
<tr id="row_0_8_1_0_4_0_" style="display:none;"><td class="entry"><span style="width:96px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1reference_1_1device_1_1detail_1_1TensorFillDiagonalFunc_1_1Params.html" target="_self">Params</a></td><td class="desc">Parameters structure </td></tr>
<tr id="row_0_8_1_0_5_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span id="arr_0_8_1_0_5_" class="arrow" onclick="toggleFolder('0_8_1_0_5_')">&#9658;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1reference_1_1device_1_1detail_1_1TensorFillLinearFunc.html" target="_self">TensorFillLinearFunc</a></td><td class="desc">Computes a random Gaussian distribution </td></tr>
<tr id="row_0_8_1_0_5_0_" style="display:none;"><td class="entry"><span style="width:96px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1reference_1_1device_1_1detail_1_1TensorFillLinearFunc_1_1Params.html" target="_self">Params</a></td><td class="desc">Parameters structure </td></tr>
<tr id="row_0_8_1_0_6_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span id="arr_0_8_1_0_6_" class="arrow" onclick="toggleFolder('0_8_1_0_6_')">&#9658;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1reference_1_1device_1_1detail_1_1TensorFillRandomGaussianFunc.html" target="_self">TensorFillRandomGaussianFunc</a></td><td class="desc">Computes a random Gaussian distribution </td></tr>
<tr id="row_0_8_1_0_6_0_" style="display:none;"><td class="entry"><span style="width:96px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1reference_1_1device_1_1detail_1_1TensorFillRandomGaussianFunc_1_1Params.html" target="_self">Params</a></td><td class="desc">Parameters structure </td></tr>
<tr id="row_0_8_1_0_7_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span id="arr_0_8_1_0_7_" class="arrow" onclick="toggleFolder('0_8_1_0_7_')">&#9658;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1reference_1_1device_1_1detail_1_1TensorFillRandomUniformFunc.html" target="_self">TensorFillRandomUniformFunc</a></td><td class="desc">Computes a random Gaussian distribution </td></tr>
<tr id="row_0_8_1_0_7_0_" style="display:none;"><td class="entry"><span style="width:96px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1reference_1_1device_1_1detail_1_1TensorFillRandomUniformFunc_1_1Params.html" target="_self">Params</a></td><td class="desc">Parameters structure </td></tr>
<tr id="row_0_8_1_0_8_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span id="arr_0_8_1_0_8_" class="arrow" onclick="toggleFolder('0_8_1_0_8_')">&#9658;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1reference_1_1device_1_1detail_1_1TensorUpdateDiagonalFunc.html" target="_self">TensorUpdateDiagonalFunc</a></td><td class="desc">Computes a random Gaussian distribution </td></tr>
<tr id="row_0_8_1_0_8_0_" style="display:none;"><td class="entry"><span style="width:96px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1reference_1_1device_1_1detail_1_1TensorUpdateDiagonalFunc_1_1Params.html" target="_self">Params</a></td><td class="desc">Parameters structure </td></tr>
<tr id="row_0_8_1_0_9_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span id="arr_0_8_1_0_9_" class="arrow" onclick="toggleFolder('0_8_1_0_9_')">&#9658;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1reference_1_1device_1_1detail_1_1TensorUpdateOffDiagonalFunc.html" target="_self">TensorUpdateOffDiagonalFunc</a></td><td class="desc">Computes a random Gaussian distribution </td></tr>
<tr id="row_0_8_1_0_9_0_" style="display:none;"><td class="entry"><span style="width:96px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1reference_1_1device_1_1detail_1_1TensorUpdateOffDiagonalFunc_1_1Params.html" target="_self">Params</a></td><td class="desc">Parameters structure </td></tr>
<tr id="row_0_8_1_1_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span id="arr_0_8_1_1_" class="arrow" onclick="toggleFolder('0_8_1_1_')">&#9658;</span><span class="icona"><span class="icon">N</span></span><a class="el" href="namespacecutlass_1_1reference_1_1device_1_1kernel.html" target="_self">kernel</a></td><td class="desc"></td></tr>
<tr id="row_0_8_1_1_0_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span id="arr_0_8_1_1_0_" class="arrow" onclick="toggleFolder('0_8_1_1_0_')">&#9658;</span><span class="icona"><span class="icon">N</span></span><a class="el" href="namespacecutlass_1_1reference_1_1device_1_1kernel_1_1detail.html" target="_self">detail</a></td><td class="desc">Defines several helpers </td></tr>
<tr id="row_0_8_1_1_0_0_" style="display:none;"><td class="entry"><span style="width:96px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1reference_1_1device_1_1kernel_1_1detail_1_1TensorForEachHelper.html" target="_self">TensorForEachHelper</a></td><td class="desc">Helper to perform for-each operation </td></tr>
<tr id="row_0_8_1_1_0_1_" style="display:none;"><td class="entry"><span style="width:96px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1reference_1_1device_1_1kernel_1_1detail_1_1TensorForEachHelper_3_01Func_00_01Rank_00_010_01_4.html" target="_self">TensorForEachHelper&lt; Func, Rank, 0 &gt;</a></td><td class="desc">Helper to perform for-each operation </td></tr>
<tr id="row_0_8_1_2_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span id="arr_0_8_1_2_" class="arrow" onclick="toggleFolder('0_8_1_2_')">&#9658;</span><span class="icona"><span class="icon">N</span></span><a class="el" href="namespacecutlass_1_1reference_1_1device_1_1thread.html" target="_self">thread</a></td><td class="desc"></td></tr>
<tr id="row_0_8_1_2_0_" style="display:none;"><td class="entry"><span style="width:80px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1reference_1_1device_1_1thread_1_1Gemm.html" target="_self">Gemm</a></td><td class="desc">Thread-level blocked general matrix product </td></tr>
<tr id="row_0_8_1_3_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1reference_1_1device_1_1BlockForEach.html" target="_self">BlockForEach</a></td><td class="desc"></td></tr>
<tr id="row_0_8_1_4_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1reference_1_1device_1_1Gemm.html" target="_self">Gemm</a></td><td class="desc"></td></tr>
<tr id="row_0_8_1_5_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1reference_1_1device_1_1Gemm_3_01ElementA_00_01LayoutA_00_01ElementB_00_01Layout4e016ab7cfc644acd7cb4ae770339773.html" target="_self">Gemm&lt; ElementA, LayoutA, ElementB, LayoutB, ElementC, LayoutC, ScalarType, AccumulatorType, arch::OpMultiplyAdd &gt;</a></td><td class="desc">Partial specialization for multiply-add </td></tr>
<tr id="row_0_8_1_6_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1reference_1_1device_1_1Gemm_3_01ElementA_00_01LayoutA_00_01ElementB_00_01Layout30b72addd464a2ca4a26785cbfd77a8e.html" target="_self">Gemm&lt; ElementA, LayoutA, ElementB, LayoutB, ElementC, LayoutC, ScalarType, AccumulatorType, arch::OpMultiplyAddSaturate &gt;</a></td><td class="desc">Partial specialization for multiply-add-saturate </td></tr>
<tr id="row_0_8_1_7_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1reference_1_1device_1_1Gemm_3_01ElementA_00_01LayoutA_00_01ElementB_00_01Layout660562b232f408218828ca5915b7e73a.html" target="_self">Gemm&lt; ElementA, LayoutA, ElementB, LayoutB, ElementC, LayoutC, ScalarType, AccumulatorType, arch::OpXorPopc &gt;</a></td><td class="desc">Parital specialization for XOR-popc </td></tr>
<tr id="row_0_8_1_8_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1reference_1_1device_1_1TensorDiagonalForEach.html" target="_self">TensorDiagonalForEach</a></td><td class="desc">Launches a kernel calling a functor for each element along a tensor's diagonal </td></tr>
<tr id="row_0_8_1_9_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1reference_1_1device_1_1TensorForEach.html" target="_self">TensorForEach</a></td><td class="desc">Launches a kernel calling a functor for each element in a tensor's index space </td></tr>
<tr id="row_0_8_2_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span id="arr_0_8_2_" class="arrow" onclick="toggleFolder('0_8_2_')">&#9658;</span><span class="icona"><span class="icon">N</span></span><a class="el" href="namespacecutlass_1_1reference_1_1host.html" target="_self">host</a></td><td class="desc"></td></tr>
<tr id="row_0_8_2_0_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span id="arr_0_8_2_0_" class="arrow" onclick="toggleFolder('0_8_2_0_')">&#9658;</span><span class="icona"><span class="icon">N</span></span><a class="el" href="namespacecutlass_1_1reference_1_1host_1_1detail.html" target="_self">detail</a></td><td class="desc">Defines several helpers </td></tr>
<tr id="row_0_8_2_0_0_" style="display:none;"><td class="entry"><span style="width:80px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1reference_1_1host_1_1detail_1_1RandomGaussianFunc.html" target="_self">RandomGaussianFunc</a></td><td class="desc"></td></tr>
<tr id="row_0_8_2_0_1_" style="display:none;"><td class="entry"><span style="width:80px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1reference_1_1host_1_1detail_1_1RandomGaussianFunc_3_01complex_3_01Element_01_4_01_4.html" target="_self">RandomGaussianFunc&lt; complex&lt; Element &gt; &gt;</a></td><td class="desc">Partial specialization for initializing a complex value </td></tr>
<tr id="row_0_8_2_0_2_" style="display:none;"><td class="entry"><span style="width:80px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1reference_1_1host_1_1detail_1_1RandomUniformFunc.html" target="_self">RandomUniformFunc</a></td><td class="desc"></td></tr>
<tr id="row_0_8_2_0_3_" style="display:none;"><td class="entry"><span style="width:80px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1reference_1_1host_1_1detail_1_1RandomUniformFunc_3_01complex_3_01Element_01_4_01_4.html" target="_self">RandomUniformFunc&lt; complex&lt; Element &gt; &gt;</a></td><td class="desc">Partial specialization for initializing a complex value </td></tr>
<tr id="row_0_8_2_0_4_" style="display:none;"><td class="entry"><span style="width:80px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1reference_1_1host_1_1detail_1_1TensorContainsFunc.html" target="_self">TensorContainsFunc</a></td><td class="desc">&lt; Layout function </td></tr>
<tr id="row_0_8_2_0_5_" style="display:none;"><td class="entry"><span style="width:80px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1reference_1_1host_1_1detail_1_1TensorCopyIf.html" target="_self">TensorCopyIf</a></td><td class="desc">Helper to conditionally copy between tensor views </td></tr>
<tr id="row_0_8_2_0_6_" style="display:none;"><td class="entry"><span style="width:80px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1reference_1_1host_1_1detail_1_1TensorEqualsFunc.html" target="_self">TensorEqualsFunc</a></td><td class="desc">&lt; Layout function </td></tr>
<tr id="row_0_8_2_0_7_" style="display:none;"><td class="entry"><span style="width:80px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1reference_1_1host_1_1detail_1_1TensorFillDiagonalFunc.html" target="_self">TensorFillDiagonalFunc</a></td><td class="desc">&lt; Layout function </td></tr>
<tr id="row_0_8_2_0_8_" style="display:none;"><td class="entry"><span style="width:80px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1reference_1_1host_1_1detail_1_1TensorFillFunc.html" target="_self">TensorFillFunc</a></td><td class="desc">&lt; Layout function </td></tr>
<tr id="row_0_8_2_0_9_" style="display:none;"><td class="entry"><span style="width:80px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1reference_1_1host_1_1detail_1_1TensorFillGaussianFunc.html" target="_self">TensorFillGaussianFunc</a></td><td class="desc">Computes a random Gaussian distribution </td></tr>
<tr id="row_0_8_2_0_10_" style="display:none;"><td class="entry"><span style="width:80px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1reference_1_1host_1_1detail_1_1TensorFillLinearFunc.html" target="_self">TensorFillLinearFunc</a></td><td class="desc">&lt; Layout function </td></tr>
<tr id="row_0_8_2_0_11_" style="display:none;"><td class="entry"><span style="width:80px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1reference_1_1host_1_1detail_1_1TensorFillRandomUniformFunc.html" target="_self">TensorFillRandomUniformFunc</a></td><td class="desc">Computes a random Gaussian distribution </td></tr>
<tr id="row_0_8_2_0_12_" style="display:none;"><td class="entry"><span style="width:80px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1reference_1_1host_1_1detail_1_1TensorForEachHelper.html" target="_self">TensorForEachHelper</a></td><td class="desc">Helper to perform for-each operation </td></tr>
<tr id="row_0_8_2_0_13_" style="display:none;"><td class="entry"><span style="width:80px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1reference_1_1host_1_1detail_1_1TensorForEachHelper_3_01Func_00_01Rank_00_010_01_4.html" target="_self">TensorForEachHelper&lt; Func, Rank, 0 &gt;</a></td><td class="desc">Helper to perform for-each operation </td></tr>
<tr id="row_0_8_2_0_14_" style="display:none;"><td class="entry"><span style="width:80px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1reference_1_1host_1_1detail_1_1TensorFuncBinaryOp.html" target="_self">TensorFuncBinaryOp</a></td><td class="desc">Helper to apply a binary operator in place </td></tr>
<tr id="row_0_8_2_0_15_" style="display:none;"><td class="entry"><span style="width:80px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1reference_1_1host_1_1detail_1_1TensorUpdateOffDiagonalFunc.html" target="_self">TensorUpdateOffDiagonalFunc</a></td><td class="desc">&lt; Layout function </td></tr>
<tr id="row_0_8_2_0_16_" style="display:none;"><td class="entry"><span style="width:80px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1reference_1_1host_1_1detail_1_1TrivialConvert.html" target="_self">TrivialConvert</a></td><td class="desc">Helper to convert between types </td></tr>
<tr id="row_0_8_2_1_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1reference_1_1host_1_1BlockForEach.html" target="_self">BlockForEach</a></td><td class="desc"></td></tr>
<tr id="row_0_8_2_2_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1reference_1_1host_1_1Gemm.html" target="_self">Gemm</a></td><td class="desc"></td></tr>
<tr id="row_0_8_2_3_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1reference_1_1host_1_1Gemm_3_01ElementA_00_01LayoutA_00_01ElementB_00_01LayoutB_193dd3a37f00deff1e5dcd7c310afb1f.html" target="_self">Gemm&lt; ElementA, LayoutA, ElementB, LayoutB, ElementC, LayoutC, ScalarType, ComputeType, arch::OpMultiplyAdd &gt;</a></td><td class="desc">Partial specialization for multiply-add </td></tr>
<tr id="row_0_8_2_4_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1reference_1_1host_1_1Gemm_3_01ElementA_00_01LayoutA_00_01ElementB_00_01LayoutB_55729eac7dbd6bf311ea36f680e83e93.html" target="_self">Gemm&lt; ElementA, LayoutA, ElementB, LayoutB, ElementC, LayoutC, ScalarType, ComputeType, arch::OpMultiplyAddSaturate &gt;</a></td><td class="desc">Partial specialization for multiply-add-saturate </td></tr>
<tr id="row_0_8_2_5_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1reference_1_1host_1_1Gemm_3_01ElementA_00_01LayoutA_00_01ElementB_00_01LayoutB_4f3f32c4b336238abfd741e87bfced46.html" target="_self">Gemm&lt; ElementA, LayoutA, ElementB, LayoutB, ElementC, LayoutC, ScalarType, ComputeType, arch::OpXorPopc &gt;</a></td><td class="desc">Parital specialization for XOR-popc </td></tr>
<tr id="row_0_9_" style="display:none;"><td class="entry"><span style="width:16px;display:inline-block;">&#160;</span><span id="arr_0_9_" class="arrow" onclick="toggleFolder('0_9_')">&#9658;</span><span class="icona"><span class="icon">N</span></span><a class="el" href="namespacecutlass_1_1thread.html" target="_self">thread</a></td><td class="desc"></td></tr>
<tr id="row_0_9_0_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1thread_1_1Matrix.html" target="_self">Matrix</a></td><td class="desc">Per-thread matrix object storing a packed matrix </td></tr>
<tr id="row_0_10_" style="display:none;"><td class="entry"><span style="width:16px;display:inline-block;">&#160;</span><span id="arr_0_10_" class="arrow" onclick="toggleFolder('0_10_')">&#9658;</span><span class="icona"><span class="icon">N</span></span><a class="el" href="namespacecutlass_1_1transform.html" target="_self">transform</a></td><td class="desc"></td></tr>
<tr id="row_0_10_0_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span id="arr_0_10_0_" class="arrow" onclick="toggleFolder('0_10_0_')">&#9658;</span><span class="icona"><span class="icon">N</span></span><a class="el" href="namespacecutlass_1_1transform_1_1thread.html" target="_self">thread</a></td><td class="desc"></td></tr>
<tr id="row_0_10_0_0_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1transform_1_1thread_1_1Transpose.html" target="_self">Transpose</a></td><td class="desc">Transforms a fragment by doing a transpose </td></tr>
<tr id="row_0_10_0_1_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1transform_1_1thread_1_1Transpose_3_01ElementCount___00_01layout_1_1PitchLinearS99f8e05faf0bb5ed48a0154afe740d81.html" target="_self">Transpose&lt; ElementCount_, layout::PitchLinearShape&lt; 4, 4 &gt;, int8_t &gt;</a></td><td class="desc">Specialization for int8_t 4x4 transpose </td></tr>
<tr id="row_0_10_1_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span id="arr_0_10_1_" class="arrow" onclick="toggleFolder('0_10_1_')">&#9658;</span><span class="icona"><span class="icon">N</span></span><a class="el" href="namespacecutlass_1_1transform_1_1threadblock.html" target="_self">threadblock</a></td><td class="desc"></td></tr>
<tr id="row_0_10_1_0_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1transform_1_1threadblock_1_1PredicatedTileAccessIterator.html" target="_self">PredicatedTileAccessIterator</a></td><td class="desc"></td></tr>
<tr id="row_0_10_1_1_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1transform_1_1threadblock_1_1PredicatedTileAccessIterator2dThreadTile.html" target="_self">PredicatedTileAccessIterator2dThreadTile</a></td><td class="desc"></td></tr>
<tr id="row_0_10_1_2_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span id="arr_0_10_1_2_" class="arrow" onclick="toggleFolder('0_10_1_2_')">&#9658;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1transform_1_1threadblock_1_1PredicatedTileAccessIterator2dThreadTile_3_01Shape__da632779aba661c0f4cfaaa78126b771.html" target="_self">PredicatedTileAccessIterator2dThreadTile&lt; Shape_, Element_, layout::ColumnMajor, AdvanceRank, ThreadMap_, AccessType_ &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_10_1_2_0_" style="display:none;"><td class="entry"><span style="width:80px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1transform_1_1threadblock_1_1PredicatedTileAccessIterator2dThreadTile_3_01Shape__18e9cf25bb3b8edfaad595241a6dc2d7.html" target="_self">Params</a></td><td class="desc">Parameters object is precomputed state and is host-constructible </td></tr>
<tr id="row_0_10_1_3_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span id="arr_0_10_1_3_" class="arrow" onclick="toggleFolder('0_10_1_3_')">&#9658;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1transform_1_1threadblock_1_1PredicatedTileAccessIterator2dThreadTile_3_01Shape__1790abaa54a01f277d75766d5882fec8.html" target="_self">PredicatedTileAccessIterator2dThreadTile&lt; Shape_, Element_, layout::PitchLinear, AdvanceRank, ThreadMap_, AccessType_ &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_10_1_3_0_" style="display:none;"><td class="entry"><span style="width:80px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1transform_1_1threadblock_1_1PredicatedTileAccessIterator2dThreadTile_3_01Shape__8ccc62d47a092afc8bee32ffe9d1e4ba.html" target="_self">Params</a></td><td class="desc">Parameters object is precomputed state and is host-constructible </td></tr>
<tr id="row_0_10_1_4_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span id="arr_0_10_1_4_" class="arrow" onclick="toggleFolder('0_10_1_4_')">&#9658;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1transform_1_1threadblock_1_1PredicatedTileAccessIterator2dThreadTile_3_01Shape__7327fa15996bcb8502cdfcc192350fe1.html" target="_self">PredicatedTileAccessIterator2dThreadTile&lt; Shape_, Element_, layout::RowMajor, AdvanceRank, ThreadMap_, AccessType_ &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_10_1_4_0_" style="display:none;"><td class="entry"><span style="width:80px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1transform_1_1threadblock_1_1PredicatedTileAccessIterator2dThreadTile_3_01Shape__a56cbccec33ee916292ad9d068474609.html" target="_self">Params</a></td><td class="desc">Parameters object is precomputed state and is host-constructible </td></tr>
<tr id="row_0_10_1_5_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span id="arr_0_10_1_5_" class="arrow" onclick="toggleFolder('0_10_1_5_')">&#9658;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1transform_1_1threadblock_1_1PredicatedTileAccessIterator_3_01Shape___00_01Elemen89c687c583745a73cb485041911a4c4e.html" target="_self">PredicatedTileAccessIterator&lt; Shape_, Element_, layout::ColumnMajor, AdvanceRank, ThreadMap_, AccessType_ &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_10_1_5_0_" style="display:none;"><td class="entry"><span style="width:80px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1transform_1_1threadblock_1_1PredicatedTileAccessIterator_3_01Shape___00_01Elemenc07b5ec72f83e782121ac629288d61fe.html" target="_self">Params</a></td><td class="desc">Parameters object is precomputed state and is host-constructible </td></tr>
<tr id="row_0_10_1_6_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span id="arr_0_10_1_6_" class="arrow" onclick="toggleFolder('0_10_1_6_')">&#9658;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1transform_1_1threadblock_1_1PredicatedTileAccessIterator_3_01Shape___00_01Elemenab63a1e105bf37f6371516cb9e2c5a7a.html" target="_self">PredicatedTileAccessIterator&lt; Shape_, Element_, layout::ColumnMajorInterleaved&lt; InterleavedK &gt;, AdvanceRank, ThreadMap_, AccessType_ &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_10_1_6_0_" style="display:none;"><td class="entry"><span style="width:80px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1transform_1_1threadblock_1_1PredicatedTileAccessIterator_3_01Shape___00_01Elemena9b06926a275b569ee9f7f142604b997.html" target="_self">Params</a></td><td class="desc">Parameters object is precomputed state and is host-constructible </td></tr>
<tr id="row_0_10_1_7_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span id="arr_0_10_1_7_" class="arrow" onclick="toggleFolder('0_10_1_7_')">&#9658;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1transform_1_1threadblock_1_1PredicatedTileAccessIterator_3_01Shape___00_01Elemen784a0e9da3f55064c47e5613791f51f7.html" target="_self">PredicatedTileAccessIterator&lt; Shape_, Element_, layout::PitchLinear, AdvanceRank, ThreadMap_, AccessType_ &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_10_1_7_0_" style="display:none;"><td class="entry"><span style="width:80px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1transform_1_1threadblock_1_1PredicatedTileAccessIterator_3_01Shape___00_01Elemen41e459f664d17473570cf22fb616845f.html" target="_self">Params</a></td><td class="desc">Parameters object is precomputed state and is host-constructible </td></tr>
<tr id="row_0_10_1_8_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span id="arr_0_10_1_8_" class="arrow" onclick="toggleFolder('0_10_1_8_')">&#9658;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1transform_1_1threadblock_1_1PredicatedTileAccessIterator_3_01Shape___00_01Elemen9838736ad62fae54213fbaf722a989ab.html" target="_self">PredicatedTileAccessIterator&lt; Shape_, Element_, layout::RowMajor, AdvanceRank, ThreadMap_, AccessType_ &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_10_1_8_0_" style="display:none;"><td class="entry"><span style="width:80px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1transform_1_1threadblock_1_1PredicatedTileAccessIterator_3_01Shape___00_01Elemen44ce348364e78f5a56fa0c2cef6af930.html" target="_self">Params</a></td><td class="desc">Parameters object is precomputed state and is host-constructible </td></tr>
<tr id="row_0_10_1_9_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span id="arr_0_10_1_9_" class="arrow" onclick="toggleFolder('0_10_1_9_')">&#9658;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1transform_1_1threadblock_1_1PredicatedTileAccessIterator_3_01Shape___00_01Elemen809793e785fb4211888c6b4e5dcfcb39.html" target="_self">PredicatedTileAccessIterator&lt; Shape_, Element_, layout::RowMajorInterleaved&lt; InterleavedK &gt;, AdvanceRank, ThreadMap_, AccessType_ &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_10_1_9_0_" style="display:none;"><td class="entry"><span style="width:80px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1transform_1_1threadblock_1_1PredicatedTileAccessIterator_3_01Shape___00_01Elemen058417e2cdd86f3cd6ad5458581571c8.html" target="_self">Params</a></td><td class="desc">Parameters object is precomputed state and is host-constructible </td></tr>
<tr id="row_0_10_1_10_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1transform_1_1threadblock_1_1PredicatedTileIterator.html" target="_self">PredicatedTileIterator</a></td><td class="desc"></td></tr>
<tr id="row_0_10_1_11_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1transform_1_1threadblock_1_1PredicatedTileIterator2dThreadTile.html" target="_self">PredicatedTileIterator2dThreadTile</a></td><td class="desc"></td></tr>
<tr id="row_0_10_1_12_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span id="arr_0_10_1_12_" class="arrow" onclick="toggleFolder('0_10_1_12_')">&#9658;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1transform_1_1threadblock_1_1PredicatedTileIterator2dThreadTile_3_01Shape___00_0165b39a630d10785a3558406f9adb99b9.html" target="_self">PredicatedTileIterator2dThreadTile&lt; Shape_, Element_, layout::ColumnMajor, AdvanceRank, ThreadMap_, Transpose_ &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_10_1_12_0_" style="display:none;"><td class="entry"><span style="width:80px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1transform_1_1threadblock_1_1PredicatedTileIterator2dThreadTile_3_01Shape___00_01e11ed7192af5d7ad1bce5641fa13112e.html" target="_self">Params</a></td><td class="desc">Parameters object is precomputed state and is host-constructible </td></tr>
<tr id="row_0_10_1_13_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span id="arr_0_10_1_13_" class="arrow" onclick="toggleFolder('0_10_1_13_')">&#9658;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1transform_1_1threadblock_1_1PredicatedTileIterator2dThreadTile_3_01Shape___00_017a517f3c73efd795ab05059cc9b111e1.html" target="_self">PredicatedTileIterator2dThreadTile&lt; Shape_, Element_, layout::PitchLinear, AdvanceRank, ThreadMap_, Transpose_ &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_10_1_13_0_" style="display:none;"><td class="entry"><span style="width:80px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1transform_1_1threadblock_1_1PredicatedTileIterator2dThreadTile_3_01Shape___00_0b878062cc0cd214bf7e17d74ff17e246.html" target="_self">AccessType</a></td><td class="desc"></td></tr>
<tr id="row_0_10_1_13_1_" style="display:none;"><td class="entry"><span style="width:80px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1transform_1_1threadblock_1_1PredicatedTileIterator2dThreadTile_3_01Shape___00_0145ef045e8f7d57dc718098adcb00cf3d.html" target="_self">Params</a></td><td class="desc">Parameters object is precomputed state and is host-constructible </td></tr>
<tr id="row_0_10_1_14_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span id="arr_0_10_1_14_" class="arrow" onclick="toggleFolder('0_10_1_14_')">&#9658;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1transform_1_1threadblock_1_1PredicatedTileIterator2dThreadTile_3_01Shape___00_013671177d6219bfeb0e1b4dc4c1b5bf11.html" target="_self">PredicatedTileIterator2dThreadTile&lt; Shape_, Element_, layout::RowMajor, AdvanceRank, ThreadMap_, Transpose_ &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_10_1_14_0_" style="display:none;"><td class="entry"><span style="width:80px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1transform_1_1threadblock_1_1PredicatedTileIterator2dThreadTile_3_01Shape___00_0102e766863c6ac9ec2063a02c4803eecb.html" target="_self">Params</a></td><td class="desc">Parameters object is precomputed state and is host-constructible </td></tr>
<tr id="row_0_10_1_15_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span id="arr_0_10_1_15_" class="arrow" onclick="toggleFolder('0_10_1_15_')">&#9658;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1transform_1_1threadblock_1_1PredicatedTileIterator_3_01Shape___00_01Element___0068b3e874b5d93d11f0fa902c7f1d11d9.html" target="_self">PredicatedTileIterator&lt; Shape_, Element_, layout::ColumnMajor, AdvanceRank, ThreadMap_, AccessSize &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_10_1_15_0_" style="display:none;"><td class="entry"><span style="width:80px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1transform_1_1threadblock_1_1PredicatedTileIterator_3_01Shape___00_01Element___00a6b756b1bcfbb35fe4a3e68ff074e380.html" target="_self">Params</a></td><td class="desc">Parameters object is precomputed state and is host-constructible </td></tr>
<tr id="row_0_10_1_16_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span id="arr_0_10_1_16_" class="arrow" onclick="toggleFolder('0_10_1_16_')">&#9658;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1transform_1_1threadblock_1_1PredicatedTileIterator_3_01Shape___00_01Element___00f6b3a9dfab5e7c72d5233f7e5e6e3b9b.html" target="_self">PredicatedTileIterator&lt; Shape_, Element_, layout::ColumnMajorInterleaved&lt; InterleavedK &gt;, AdvanceRank, ThreadMap_, AccessSize &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_10_1_16_0_" style="display:none;"><td class="entry"><span style="width:80px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1transform_1_1threadblock_1_1PredicatedTileIterator_3_01Shape___00_01Element___00ebd1a63351e1085d0b718582ec7b06c8.html" target="_self">Params</a></td><td class="desc">Parameters object is precomputed state and is host-constructible </td></tr>
<tr id="row_0_10_1_17_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span id="arr_0_10_1_17_" class="arrow" onclick="toggleFolder('0_10_1_17_')">&#9658;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1transform_1_1threadblock_1_1PredicatedTileIterator_3_01Shape___00_01Element___00e7c2c404e7aedfe60ad56bb5571306a1.html" target="_self">PredicatedTileIterator&lt; Shape_, Element_, layout::PitchLinear, AdvanceRank, ThreadMap_, AccessSize &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_10_1_17_0_" style="display:none;"><td class="entry"><span style="width:80px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1transform_1_1threadblock_1_1PredicatedTileIterator_3_01Shape___00_01Element___006a5f2f7a8271031e6cdc5daa5441f2af.html" target="_self">Params</a></td><td class="desc">Parameters object is precomputed state and is host-constructible </td></tr>
<tr id="row_0_10_1_18_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span id="arr_0_10_1_18_" class="arrow" onclick="toggleFolder('0_10_1_18_')">&#9658;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1transform_1_1threadblock_1_1PredicatedTileIterator_3_01Shape___00_01Element___0041ea81994f8af0d4d071fdb9e66b5ff0.html" target="_self">PredicatedTileIterator&lt; Shape_, Element_, layout::RowMajor, AdvanceRank, ThreadMap_, AccessSize &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_10_1_18_0_" style="display:none;"><td class="entry"><span style="width:80px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1transform_1_1threadblock_1_1PredicatedTileIterator_3_01Shape___00_01Element___004d0f9b5e19c29acc17bcdc360dafebbd.html" target="_self">Params</a></td><td class="desc">Parameters object is precomputed state and is host-constructible </td></tr>
<tr id="row_0_10_1_19_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span id="arr_0_10_1_19_" class="arrow" onclick="toggleFolder('0_10_1_19_')">&#9658;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1transform_1_1threadblock_1_1PredicatedTileIterator_3_01Shape___00_01Element___00d670f969180a8d182dffb356ebcc957e.html" target="_self">PredicatedTileIterator&lt; Shape_, Element_, layout::RowMajorInterleaved&lt; InterleavedK &gt;, AdvanceRank, ThreadMap_, AccessSize &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_10_1_19_0_" style="display:none;"><td class="entry"><span style="width:80px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1transform_1_1threadblock_1_1PredicatedTileIterator_3_01Shape___00_01Element___009fd89f6dad84238fd7d63df0a0c0364f.html" target="_self">Params</a></td><td class="desc">Parameters object is precomputed state and is host-constructible </td></tr>
<tr id="row_0_10_1_20_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1transform_1_1threadblock_1_1RegularTileAccessIterator.html" target="_self">RegularTileAccessIterator</a></td><td class="desc"></td></tr>
<tr id="row_0_10_1_21_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1transform_1_1threadblock_1_1RegularTileAccessIterator_3_01Shape___00_01Element__eb7d20f8b9d69e0ae5e7ef51dc480867.html" target="_self">RegularTileAccessIterator&lt; Shape_, Element_, layout::ColumnMajor, AdvanceRank, ThreadMap_, Alignment &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_10_1_22_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1transform_1_1threadblock_1_1RegularTileAccessIterator_3_01Shape___00_01Element__2c1476eaf582bfe972793e17babfe985.html" target="_self">RegularTileAccessIterator&lt; Shape_, Element_, layout::ColumnMajorTensorOpMultiplicandCongruous&lt; sizeof_bits&lt; Element_ &gt;::value, int(128/sizeof(Element_))&gt;, AdvanceRank, ThreadMap_, Alignment &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_10_1_23_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1transform_1_1threadblock_1_1RegularTileAccessIterator_3_01Shape___00_01Element__a3c11cf1f00ef7a1efb8389ac6e4c6e0.html" target="_self">RegularTileAccessIterator&lt; Shape_, Element_, layout::ColumnMajorTensorOpMultiplicandCrosswise&lt; sizeof_bits&lt; Element_ &gt;::value, Crosswise &gt;, AdvanceRank, ThreadMap_, Alignment &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_10_1_24_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1transform_1_1threadblock_1_1RegularTileAccessIterator_3_01Shape___00_01Element__0855e9d9ab619202d2397180c1e4c4a5.html" target="_self">RegularTileAccessIterator&lt; Shape_, Element_, layout::PitchLinear, AdvanceRank, ThreadMap_, Alignment &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_10_1_25_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1transform_1_1threadblock_1_1RegularTileAccessIterator_3_01Shape___00_01Element__f04332958a49a47d6fb2b25201764630.html" target="_self">RegularTileAccessIterator&lt; Shape_, Element_, layout::RowMajor, AdvanceRank, ThreadMap_, Alignment &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_10_1_26_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1transform_1_1threadblock_1_1RegularTileAccessIterator_3_01Shape___00_01Element__6baada077236f1a368c61c5e11b45b72.html" target="_self">RegularTileAccessIterator&lt; Shape_, Element_, layout::RowMajorTensorOpMultiplicandCongruous&lt; sizeof_bits&lt; Element_ &gt;::value, int(128/sizeof(Element_))&gt;, AdvanceRank, ThreadMap_, Alignment &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_10_1_27_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1transform_1_1threadblock_1_1RegularTileAccessIterator_3_01Shape___00_01Element__0184b7188941788a96624510a4b2f876.html" target="_self">RegularTileAccessIterator&lt; Shape_, Element_, layout::RowMajorTensorOpMultiplicandCrosswise&lt; sizeof_bits&lt; Element_ &gt;::value, Crosswise &gt;, AdvanceRank, ThreadMap_, Alignment &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_10_1_28_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span id="arr_0_10_1_28_" class="arrow" onclick="toggleFolder('0_10_1_28_')">&#9658;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1transform_1_1threadblock_1_1RegularTileAccessIterator_3_01Shape___00_01Element__ebf4714349612673e8b6609b763eeb6f.html" target="_self">RegularTileAccessIterator&lt; Shape_, Element_, layout::TensorOpMultiplicandCongruous&lt; sizeof_bits&lt; Element_ &gt;::value, int(128/sizeof(Element_))&gt;, AdvanceRank, ThreadMap_, Alignment &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_10_1_28_0_" style="display:none;"><td class="entry"><span style="width:80px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1transform_1_1threadblock_1_1RegularTileAccessIterator_3_01Shape___00_01Element_0a9491607d11be8e1780e79ad711aa42.html" target="_self">Detail</a></td><td class="desc">Internal details made public to facilitate introspection </td></tr>
<tr id="row_0_10_1_29_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span id="arr_0_10_1_29_" class="arrow" onclick="toggleFolder('0_10_1_29_')">&#9658;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1transform_1_1threadblock_1_1RegularTileAccessIterator_3_01Shape___00_01Element__e9a9e0f4286f652f55eb9b863b21effe.html" target="_self">RegularTileAccessIterator&lt; Shape_, Element_, layout::TensorOpMultiplicandCrosswise&lt; sizeof_bits&lt; Element_ &gt;::value, Crosswise &gt;, AdvanceRank, ThreadMap_, Alignment &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_10_1_29_0_" style="display:none;"><td class="entry"><span style="width:80px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1transform_1_1threadblock_1_1RegularTileAccessIterator_3_01Shape___00_01Element_3be8b96d170d886f39b6b30acab65e7a.html" target="_self">Detail</a></td><td class="desc">Internal details made public to facilitate introspection </td></tr>
<tr id="row_0_10_1_30_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1transform_1_1threadblock_1_1RegularTileIterator.html" target="_self">RegularTileIterator</a></td><td class="desc"></td></tr>
<tr id="row_0_10_1_31_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1transform_1_1threadblock_1_1RegularTileIterator2dThreadTile.html" target="_self">RegularTileIterator2dThreadTile</a></td><td class="desc"></td></tr>
<tr id="row_0_10_1_32_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1transform_1_1threadblock_1_1RegularTileIterator2dThreadTile_3_01Shape___00_01Eleb60d066756d1c18f05fceee6a27bdb8a.html" target="_self">RegularTileIterator2dThreadTile&lt; Shape_, Element_, layout::ColumnMajorInterleaved&lt; 4 &gt;, AdvanceRank, ThreadMap_, Alignment &gt;</a></td><td class="desc">Regular tile iterator specialized for interleaved layout + 2d thread-tiled threadmapping </td></tr>
<tr id="row_0_10_1_33_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1transform_1_1threadblock_1_1RegularTileIterator2dThreadTile_3_01Shape___00_01Ele76ed82829532ae1c17f4c78158f036c7.html" target="_self">RegularTileIterator2dThreadTile&lt; Shape_, Element_, layout::PitchLinear, AdvanceRank, ThreadMap_, Alignment &gt;</a></td><td class="desc">Regular tile iterator specialized for pitch-linear + 2d thread-tiled threadmapping </td></tr>
<tr id="row_0_10_1_34_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1transform_1_1threadblock_1_1RegularTileIterator2dThreadTile_3_01Shape___00_01Ele654c8f6161ae5340f040397a4e2e045c.html" target="_self">RegularTileIterator2dThreadTile&lt; Shape_, Element_, layout::RowMajorInterleaved&lt; 4 &gt;, AdvanceRank, ThreadMap_, Alignment &gt;</a></td><td class="desc">Regular tile iterator specialized for interleaved layout + 2d thread-tiled threadmapping </td></tr>
<tr id="row_0_10_1_35_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1transform_1_1threadblock_1_1RegularTileIterator_3_01Shape___00_01Element___00_011d3637dbd8bc58bcb020b51bf57fbfc0.html" target="_self">RegularTileIterator&lt; Shape_, Element_, layout::ColumnMajor, AdvanceRank, ThreadMap_, Alignment &gt;</a></td><td class="desc">Regular tile iterator specialized for pitch-linear </td></tr>
<tr id="row_0_10_1_36_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1transform_1_1threadblock_1_1RegularTileIterator_3_01Shape___00_01Element___00_017982f81d4ef592e19c8427de2ea933a3.html" target="_self">RegularTileIterator&lt; Shape_, Element_, layout::ColumnMajorTensorOpMultiplicandCongruous&lt; sizeof_bits&lt; Element_ &gt;::value, int(128/sizeof(Element_))&gt;, AdvanceRank, ThreadMap_, Alignment &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_10_1_37_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1transform_1_1threadblock_1_1RegularTileIterator_3_01Shape___00_01Element___00_010889a732373c350de9b9a9f6c13cd761.html" target="_self">RegularTileIterator&lt; Shape_, Element_, layout::ColumnMajorTensorOpMultiplicandCrosswise&lt; sizeof_bits&lt; Element_ &gt;::value, Crosswise &gt;, AdvanceRank, ThreadMap_, Alignment &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_10_1_38_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1transform_1_1threadblock_1_1RegularTileIterator_3_01Shape___00_01Element___00_01187f8574e1fe9d7d5e8fbf09bd834bf0.html" target="_self">RegularTileIterator&lt; Shape_, Element_, layout::ColumnMajorVoltaTensorOpMultiplicandBCongruous&lt; sizeof_bits&lt; Element_ &gt;::value &gt;, AdvanceRank, ThreadMap_, Alignment &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_10_1_39_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1transform_1_1threadblock_1_1RegularTileIterator_3_01Shape___00_01Element___00_01793f74bfd8f116a827948ab01a37349a.html" target="_self">RegularTileIterator&lt; Shape_, Element_, layout::ColumnMajorVoltaTensorOpMultiplicandCongruous&lt; sizeof_bits&lt; Element_ &gt;::value &gt;, AdvanceRank, ThreadMap_, Alignment &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_10_1_40_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1transform_1_1threadblock_1_1RegularTileIterator_3_01Shape___00_01Element___00_01bd31b3810c1fedf2e7e5959ff92b5d3d.html" target="_self">RegularTileIterator&lt; Shape_, Element_, layout::ColumnMajorVoltaTensorOpMultiplicandCrosswise&lt; sizeof_bits&lt; Element_ &gt;::value, Shape_::kRow &gt;, AdvanceRank, ThreadMap_, Alignment &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_10_1_41_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1transform_1_1threadblock_1_1RegularTileIterator_3_01Shape___00_01Element___00_0184a89653916f5d51ab59d1b386989a17.html" target="_self">RegularTileIterator&lt; Shape_, Element_, layout::PitchLinear, AdvanceRank, ThreadMap_, Alignment &gt;</a></td><td class="desc">Regular tile iterator specialized for pitch-linear </td></tr>
<tr id="row_0_10_1_42_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1transform_1_1threadblock_1_1RegularTileIterator_3_01Shape___00_01Element___00_0149454d361ea5885cf5166a920b5145df.html" target="_self">RegularTileIterator&lt; Shape_, Element_, layout::RowMajor, AdvanceRank, ThreadMap_, Alignment &gt;</a></td><td class="desc">Regular tile iterator specialized for pitch-linear </td></tr>
<tr id="row_0_10_1_43_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1transform_1_1threadblock_1_1RegularTileIterator_3_01Shape___00_01Element___00_01c20d35180520077a5a09b1e33543c1a5.html" target="_self">RegularTileIterator&lt; Shape_, Element_, layout::RowMajorTensorOpMultiplicandCongruous&lt; sizeof_bits&lt; Element_ &gt;::value, int(128/sizeof(Element_))&gt;, AdvanceRank, ThreadMap_, Alignment &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_10_1_44_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1transform_1_1threadblock_1_1RegularTileIterator_3_01Shape___00_01Element___00_01a31b454d9c930525c1e9ca406a514f40.html" target="_self">RegularTileIterator&lt; Shape_, Element_, layout::RowMajorTensorOpMultiplicandCrosswise&lt; sizeof_bits&lt; Element_ &gt;::value, Crosswise &gt;, AdvanceRank, ThreadMap_, Alignment &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_10_1_45_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1transform_1_1threadblock_1_1RegularTileIterator_3_01Shape___00_01Element___00_0104ad31bd559a88cc418ae1cab7492ed5.html" target="_self">RegularTileIterator&lt; Shape_, Element_, layout::RowMajorVoltaTensorOpMultiplicandBCongruous&lt; sizeof_bits&lt; Element_ &gt;::value &gt;, AdvanceRank, ThreadMap_, Alignment &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_10_1_46_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1transform_1_1threadblock_1_1RegularTileIterator_3_01Shape___00_01Element___00_01f6f6511b5033cad31083644ac69c54d8.html" target="_self">RegularTileIterator&lt; Shape_, Element_, layout::RowMajorVoltaTensorOpMultiplicandCongruous&lt; sizeof_bits&lt; Element_ &gt;::value &gt;, AdvanceRank, ThreadMap_, Alignment &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_10_1_47_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1transform_1_1threadblock_1_1RegularTileIterator_3_01Shape___00_01Element___00_01b3fa5720e807697de61b9f937b269cd0.html" target="_self">RegularTileIterator&lt; Shape_, Element_, layout::RowMajorVoltaTensorOpMultiplicandCrosswise&lt; sizeof_bits&lt; Element_ &gt;::value, Shape_::kColumn &gt;, AdvanceRank, ThreadMap_, Alignment &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_10_1_48_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span id="arr_0_10_1_48_" class="arrow" onclick="toggleFolder('0_10_1_48_')">&#9658;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1transform_1_1threadblock_1_1RegularTileIterator_3_01Shape___00_01Element___00_01efd5013a2503d6567e2bf6b40c97360c.html" target="_self">RegularTileIterator&lt; Shape_, Element_, layout::TensorOpMultiplicandCongruous&lt; sizeof_bits&lt; Element_ &gt;::value, int(128/sizeof(Element_))&gt;, AdvanceRank, ThreadMap_, Alignment &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_10_1_48_0_" style="display:none;"><td class="entry"><span style="width:80px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1transform_1_1threadblock_1_1RegularTileIterator_3_01Shape___00_01Element___00_052caec9d5bceeb59b9a13cb3338ce64d.html" target="_self">Detail</a></td><td class="desc">Internal details made public to facilitate introspection </td></tr>
<tr id="row_0_10_1_49_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span id="arr_0_10_1_49_" class="arrow" onclick="toggleFolder('0_10_1_49_')">&#9658;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1transform_1_1threadblock_1_1RegularTileIterator_3_01Shape___00_01Element___00_0197fef2242a3454a7d1cebe61aee28b43.html" target="_self">RegularTileIterator&lt; Shape_, Element_, layout::TensorOpMultiplicandCrosswise&lt; sizeof_bits&lt; Element_ &gt;::value, Crosswise &gt;, AdvanceRank, ThreadMap_, Alignment &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_10_1_49_0_" style="display:none;"><td class="entry"><span style="width:80px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1transform_1_1threadblock_1_1RegularTileIterator_3_01Shape___00_01Element___00_039093927f4b1ee61538c569bf1ae4efd.html" target="_self">Detail</a></td><td class="desc">Internal details made public to facilitate introspection </td></tr>
<tr id="row_0_10_1_50_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span id="arr_0_10_1_50_" class="arrow" onclick="toggleFolder('0_10_1_50_')">&#9658;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1transform_1_1threadblock_1_1RegularTileIterator_3_01Shape___00_01Element___00_01a75d2cd74e722d6ad6a3b41aabfd432d.html" target="_self">RegularTileIterator&lt; Shape_, Element_, layout::VoltaTensorOpMultiplicandBCongruous&lt; sizeof_bits&lt; Element_ &gt;::value &gt;, AdvanceRank, ThreadMap_, Alignment &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_10_1_50_0_" style="display:none;"><td class="entry"><span style="width:80px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1transform_1_1threadblock_1_1RegularTileIterator_3_01Shape___00_01Element___00_02d305cfb0b55c6fb236a52cf2240651e.html" target="_self">Detail</a></td><td class="desc">Internal details made public to facilitate introspection </td></tr>
<tr id="row_0_10_1_51_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span id="arr_0_10_1_51_" class="arrow" onclick="toggleFolder('0_10_1_51_')">&#9658;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1transform_1_1threadblock_1_1RegularTileIterator_3_01Shape___00_01Element___00_01f96bbeb63e6d4ce4a2551279de3a9f0e.html" target="_self">RegularTileIterator&lt; Shape_, Element_, layout::VoltaTensorOpMultiplicandCongruous&lt; sizeof_bits&lt; Element_ &gt;::value &gt;, AdvanceRank, ThreadMap_, Alignment &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_10_1_51_0_" style="display:none;"><td class="entry"><span style="width:80px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1transform_1_1threadblock_1_1RegularTileIterator_3_01Shape___00_01Element___00_032f88d1be8b209e44a4815c707ba35bb.html" target="_self">Detail</a></td><td class="desc">Internal details made public to facilitate introspection </td></tr>
<tr id="row_0_10_1_52_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span id="arr_0_10_1_52_" class="arrow" onclick="toggleFolder('0_10_1_52_')">&#9658;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1transform_1_1threadblock_1_1RegularTileIterator_3_01Shape___00_01Element___00_01dbd6b8468d5bd787308d2f615a24d123.html" target="_self">RegularTileIterator&lt; Shape_, Element_, layout::VoltaTensorOpMultiplicandCrosswise&lt; sizeof_bits&lt; Element_ &gt;::value, Shape_::kContiguous &gt;, AdvanceRank, ThreadMap_, Alignment &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_10_1_52_0_" style="display:none;"><td class="entry"><span style="width:80px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1transform_1_1threadblock_1_1RegularTileIterator_3_01Shape___00_01Element___00_0390833403016f5d817416e20828845df.html" target="_self">Detail</a></td><td class="desc">Internal details made public to facilitate introspection </td></tr>
<tr id="row_0_10_2_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1transform_1_1PitchLinear2DThreadTileStripminedThreadMap.html" target="_self">PitchLinear2DThreadTileStripminedThreadMap</a></td><td class="desc"></td></tr>
<tr id="row_0_10_3_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span id="arr_0_10_3_" class="arrow" onclick="toggleFolder('0_10_3_')">&#9658;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1transform_1_1PitchLinear2DThreadTileStripminedThreadMap_3_01Shape___00_01Thread0082c3467229b12cc9dd996283ee7160.html" target="_self">PitchLinear2DThreadTileStripminedThreadMap&lt; Shape_, Threads, cutlass::layout::PitchLinearShape&lt; 4, 4 &gt; &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_10_3_0_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1transform_1_1PitchLinear2DThreadTileStripminedThreadMap_3_01Shape___00_01Thread896c01a3c466da1bf392e0cdfced4d53.html" target="_self">Detail</a></td><td class="desc">Internal implementation details </td></tr>
<tr id="row_0_10_4_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span id="arr_0_10_4_" class="arrow" onclick="toggleFolder('0_10_4_')">&#9658;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1transform_1_1PitchLinearStripminedThreadMap.html" target="_self">PitchLinearStripminedThreadMap</a></td><td class="desc"></td></tr>
<tr id="row_0_10_4_0_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1transform_1_1PitchLinearStripminedThreadMap_1_1Detail.html" target="_self">Detail</a></td><td class="desc">Internal implementation details </td></tr>
<tr id="row_0_10_5_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1transform_1_1PitchLinearTilePolicyStripminedThreadContiguous.html" target="_self">PitchLinearTilePolicyStripminedThreadContiguous</a></td><td class="desc"></td></tr>
<tr id="row_0_10_6_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1transform_1_1PitchLinearTilePolicyStripminedThreadStrided.html" target="_self">PitchLinearTilePolicyStripminedThreadStrided</a></td><td class="desc"></td></tr>
<tr id="row_0_10_7_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span id="arr_0_10_7_" class="arrow" onclick="toggleFolder('0_10_7_')">&#9658;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1transform_1_1PitchLinearWarpRakedThreadMap.html" target="_self">PitchLinearWarpRakedThreadMap</a></td><td class="desc"></td></tr>
<tr id="row_0_10_7_0_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1transform_1_1PitchLinearWarpRakedThreadMap_1_1Detail.html" target="_self">Detail</a></td><td class="desc">Internal details made public to facilitate introspection Iterations along each dimension (concept: PitchLinearShape) </td></tr>
<tr id="row_0_10_8_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span id="arr_0_10_8_" class="arrow" onclick="toggleFolder('0_10_8_')">&#9658;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1transform_1_1PitchLinearWarpStripedThreadMap.html" target="_self">PitchLinearWarpStripedThreadMap</a></td><td class="desc"></td></tr>
<tr id="row_0_10_8_0_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1transform_1_1PitchLinearWarpStripedThreadMap_1_1Detail.html" target="_self">Detail</a></td><td class="desc">Internal details made public to facilitate introspection Iterations along each dimension (concept: PitchLinearShape) </td></tr>
<tr id="row_0_10_9_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span id="arr_0_10_9_" class="arrow" onclick="toggleFolder('0_10_9_')">&#9658;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1transform_1_1TransposePitchLinearThreadMap.html" target="_self">TransposePitchLinearThreadMap</a></td><td class="desc"></td></tr>
<tr id="row_0_10_9_0_" style="display:none;"><td class="entry"><span style="width:64px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1transform_1_1TransposePitchLinearThreadMap_1_1Detail.html" target="_self">Detail</a></td><td class="desc">Internal details made public to facilitate introspection Iterations along each dimension (concept: PitchLinearShape) </td></tr>
<tr id="row_0_10_10_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1transform_1_1TransposePitchLinearThreadMap2DThreadTile.html" target="_self">TransposePitchLinearThreadMap2DThreadTile</a></td><td class="desc">Thread Mapping a 2D threadtiled mapping as a transposed Pitchlinear2DThreadTile mapping </td></tr>
<tr id="row_0_10_11_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1transform_1_1TransposePitchLinearThreadMapSimt.html" target="_self">TransposePitchLinearThreadMapSimt</a></td><td class="desc"></td></tr>
<tr id="row_0_11_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1AlignedArray.html" target="_self">AlignedArray</a></td><td class="desc">Aligned array type </td></tr>
<tr id="row_0_12_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1AlignedBuffer.html" target="_self">AlignedBuffer</a></td><td class="desc">Modifies semantics of cutlass::Array&lt;&gt; to provide guaranteed alignment </td></tr>
<tr id="row_0_13_" style="display:none;"><td class="entry"><span style="width:16px;display:inline-block;">&#160;</span><span id="arr_0_13_" class="arrow" onclick="toggleFolder('0_13_')">&#9658;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html" target="_self">Array&lt; T, N, false &gt;</a></td><td class="desc">Statically sized array for any data type </td></tr>
<tr id="row_0_13_0_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1const__iterator.html" target="_self">const_iterator</a></td><td class="desc">Bidirectional constant iterator over elements </td></tr>
<tr id="row_0_13_1_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1const__reference.html" target="_self">const_reference</a></td><td class="desc">Reference object extracts sub-byte items </td></tr>
<tr id="row_0_13_2_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1const__reverse__iterator.html" target="_self">const_reverse_iterator</a></td><td class="desc">Bidirectional constant iterator over elements </td></tr>
<tr id="row_0_13_3_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1iterator.html" target="_self">iterator</a></td><td class="desc">Bidirectional iterator over elements </td></tr>
<tr id="row_0_13_4_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1reference.html" target="_self">reference</a></td><td class="desc">Reference object inserts or extracts sub-byte items </td></tr>
<tr id="row_0_13_5_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1reverse__iterator.html" target="_self">reverse_iterator</a></td><td class="desc">Bidirectional iterator over elements </td></tr>
<tr id="row_0_14_" style="display:none;"><td class="entry"><span style="width:16px;display:inline-block;">&#160;</span><span id="arr_0_14_" class="arrow" onclick="toggleFolder('0_14_')">&#9658;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1Array_3_01T_00_01N_00_01true_01_4.html" target="_self">Array&lt; T, N, true &gt;</a></td><td class="desc">Statically sized array for any data type </td></tr>
<tr id="row_0_14_0_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1Array_3_01T_00_01N_00_01true_01_4_1_1const__iterator.html" target="_self">const_iterator</a></td><td class="desc">Bidirectional constant iterator over elements </td></tr>
<tr id="row_0_14_1_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1Array_3_01T_00_01N_00_01true_01_4_1_1const__reverse__iterator.html" target="_self">const_reverse_iterator</a></td><td class="desc">Bidirectional constant iterator over elements </td></tr>
<tr id="row_0_14_2_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1Array_3_01T_00_01N_00_01true_01_4_1_1iterator.html" target="_self">iterator</a></td><td class="desc">Bidirectional iterator over elements </td></tr>
<tr id="row_0_14_3_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1Array_3_01T_00_01N_00_01true_01_4_1_1reverse__iterator.html" target="_self">reverse_iterator</a></td><td class="desc">Bidirectional iterator over elements </td></tr>
<tr id="row_0_15_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1CommandLine.html" target="_self">CommandLine</a></td><td class="desc"></td></tr>
<tr id="row_0_16_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1complex.html" target="_self">complex</a></td><td class="desc"></td></tr>
<tr id="row_0_17_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1ConstSubbyteReference.html" target="_self">ConstSubbyteReference</a></td><td class="desc"></td></tr>
<tr id="row_0_18_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1Coord.html" target="_self">Coord</a></td><td class="desc">Statically-sized array specifying Coords within a tensor </td></tr>
<tr id="row_0_19_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1cuda__exception.html" target="_self">cuda_exception</a></td><td class="desc">C++ exception wrapper for CUDA <code>cudaError_t</code> </td></tr>
<tr id="row_0_20_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1Distribution.html" target="_self">Distribution</a></td><td class="desc"><a class="el" href="structcutlass_1_1Distribution.html" title="Distribution type. ">Distribution</a> type </td></tr>
<tr id="row_0_21_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1divide__assert.html" target="_self">divide_assert</a></td><td class="desc"></td></tr>
<tr id="row_0_22_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1divides.html" target="_self">divides</a></td><td class="desc"></td></tr>
<tr id="row_0_23_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1divides_3_01Array_3_01half__t_00_01N_01_4_01_4.html" target="_self">divides&lt; Array&lt; half_t, N &gt; &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_24_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1divides_3_01Array_3_01T_00_01N_01_4_01_4.html" target="_self">divides&lt; Array&lt; T, N &gt; &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_25_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1FloatType.html" target="_self">FloatType</a></td><td class="desc">Defines a floating-point type based on the number of exponent and mantissa bits </td></tr>
<tr id="row_0_26_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1FloatType_3_0111_00_0152_01_4.html" target="_self">FloatType&lt; 11, 52 &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_27_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1FloatType_3_015_00_0110_01_4.html" target="_self">FloatType&lt; 5, 10 &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_28_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1FloatType_3_018_00_0123_01_4.html" target="_self">FloatType&lt; 8, 23 &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_29_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1half__t.html" target="_self">half_t</a></td><td class="desc">IEEE half-precision floating-point type </td></tr>
<tr id="row_0_30_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1HostTensor.html" target="_self">HostTensor</a></td><td class="desc">Host tensor </td></tr>
<tr id="row_0_31_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1IdentityTensorLayout.html" target="_self">IdentityTensorLayout</a></td><td class="desc"></td></tr>
<tr id="row_0_32_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1integer__subbyte.html" target="_self">integer_subbyte</a></td><td class="desc">4-bit signed integer type </td></tr>
<tr id="row_0_33_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1IntegerType.html" target="_self">IntegerType</a></td><td class="desc">Defines integers based on size and whether they are signed </td></tr>
<tr id="row_0_34_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1IntegerType_3_011_00_01false_01_4.html" target="_self">IntegerType&lt; 1, false &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_35_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1IntegerType_3_011_00_01true_01_4.html" target="_self">IntegerType&lt; 1, true &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_36_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1IntegerType_3_0116_00_01false_01_4.html" target="_self">IntegerType&lt; 16, false &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_37_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1IntegerType_3_0116_00_01true_01_4.html" target="_self">IntegerType&lt; 16, true &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_38_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1IntegerType_3_0132_00_01false_01_4.html" target="_self">IntegerType&lt; 32, false &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_39_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1IntegerType_3_0132_00_01true_01_4.html" target="_self">IntegerType&lt; 32, true &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_40_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1IntegerType_3_014_00_01false_01_4.html" target="_self">IntegerType&lt; 4, false &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_41_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1IntegerType_3_014_00_01true_01_4.html" target="_self">IntegerType&lt; 4, true &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_42_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1IntegerType_3_0164_00_01false_01_4.html" target="_self">IntegerType&lt; 64, false &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_43_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1IntegerType_3_0164_00_01true_01_4.html" target="_self">IntegerType&lt; 64, true &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_44_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1IntegerType_3_018_00_01false_01_4.html" target="_self">IntegerType&lt; 8, false &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_45_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1IntegerType_3_018_00_01true_01_4.html" target="_self">IntegerType&lt; 8, true &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_46_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1is__pow2.html" target="_self">is_pow2</a></td><td class="desc"></td></tr>
<tr id="row_0_47_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1KernelLaunchConfiguration.html" target="_self">KernelLaunchConfiguration</a></td><td class="desc">Structure containing the basic launch configuration of a CUDA kernel </td></tr>
<tr id="row_0_48_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1log2__down.html" target="_self">log2_down</a></td><td class="desc"></td></tr>
<tr id="row_0_49_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1log2__down_3_01N_00_011_00_01Count_01_4.html" target="_self">log2_down&lt; N, 1, Count &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_50_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1log2__up.html" target="_self">log2_up</a></td><td class="desc"></td></tr>
<tr id="row_0_51_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1log2__up_3_01N_00_011_00_01Count_01_4.html" target="_self">log2_up&lt; N, 1, Count &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_52_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1MatrixCoord.html" target="_self">MatrixCoord</a></td><td class="desc"></td></tr>
<tr id="row_0_53_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1MatrixShape.html" target="_self">MatrixShape</a></td><td class="desc">Describes the size of a matrix tile </td></tr>
<tr id="row_0_54_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1Max.html" target="_self">Max</a></td><td class="desc"></td></tr>
<tr id="row_0_55_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1maximum.html" target="_self">maximum</a></td><td class="desc"></td></tr>
<tr id="row_0_56_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1maximum_3_01Array_3_01T_00_01N_01_4_01_4.html" target="_self">maximum&lt; Array&lt; T, N &gt; &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_57_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1maximum_3_01float_01_4.html" target="_self">maximum&lt; float &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_58_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1Min.html" target="_self">Min</a></td><td class="desc"></td></tr>
<tr id="row_0_59_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1minimum.html" target="_self">minimum</a></td><td class="desc"></td></tr>
<tr id="row_0_60_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1minimum_3_01Array_3_01T_00_01N_01_4_01_4.html" target="_self">minimum&lt; Array&lt; T, N &gt; &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_61_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1minimum_3_01float_01_4.html" target="_self">minimum&lt; float &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_62_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1minus.html" target="_self">minus</a></td><td class="desc"></td></tr>
<tr id="row_0_63_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1minus_3_01Array_3_01half__t_00_01N_01_4_01_4.html" target="_self">minus&lt; Array&lt; half_t, N &gt; &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_64_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1minus_3_01Array_3_01T_00_01N_01_4_01_4.html" target="_self">minus&lt; Array&lt; T, N &gt; &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_65_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1multiplies.html" target="_self">multiplies</a></td><td class="desc"></td></tr>
<tr id="row_0_66_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1multiplies_3_01Array_3_01half__t_00_01N_01_4_01_4.html" target="_self">multiplies&lt; Array&lt; half_t, N &gt; &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_67_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1multiplies_3_01Array_3_01T_00_01N_01_4_01_4.html" target="_self">multiplies&lt; Array&lt; T, N &gt; &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_68_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1multiply__add.html" target="_self">multiply_add</a></td><td class="desc">Fused multiply-add </td></tr>
<tr id="row_0_69_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1multiply__add_3_01Array_3_01half__t_00_01N_01_4_00_01Array_3_01half__t_00_01N_01adaeadb27c0e4439444709c0eb30963.html" target="_self">multiply_add&lt; Array&lt; half_t, N &gt;, Array&lt; half_t, N &gt;, Array&lt; half_t, N &gt; &gt;</a></td><td class="desc">Fused multiply-add </td></tr>
<tr id="row_0_70_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1multiply__add_3_01Array_3_01T_00_01N_01_4_00_01Array_3_01T_00_01N_01_4_00_01Array_3_01T_00_01N_01_4_01_4.html" target="_self">multiply_add&lt; Array&lt; T, N &gt;, Array&lt; T, N &gt;, Array&lt; T, N &gt; &gt;</a></td><td class="desc">Fused multiply-add </td></tr>
<tr id="row_0_71_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1multiply__add_3_01complex_3_01T_01_4_00_01complex_3_01T_01_4_00_01complex_3_01T_01_4_01_4.html" target="_self">multiply_add&lt; complex&lt; T &gt;, complex&lt; T &gt;, complex&lt; T &gt; &gt;</a></td><td class="desc">Fused multiply-add </td></tr>
<tr id="row_0_72_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1multiply__add_3_01complex_3_01T_01_4_00_01T_00_01complex_3_01T_01_4_01_4.html" target="_self">multiply_add&lt; complex&lt; T &gt;, T, complex&lt; T &gt; &gt;</a></td><td class="desc">Fused multiply-add </td></tr>
<tr id="row_0_73_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1multiply__add_3_01T_00_01complex_3_01T_01_4_00_01complex_3_01T_01_4_01_4.html" target="_self">multiply_add&lt; T, complex&lt; T &gt;, complex&lt; T &gt; &gt;</a></td><td class="desc">Fused multiply-add </td></tr>
<tr id="row_0_74_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1negate.html" target="_self">negate</a></td><td class="desc"></td></tr>
<tr id="row_0_75_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1negate_3_01Array_3_01half__t_00_01N_01_4_01_4.html" target="_self">negate&lt; Array&lt; half_t, N &gt; &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_76_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1negate_3_01Array_3_01T_00_01N_01_4_01_4.html" target="_self">negate&lt; Array&lt; T, N &gt; &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_77_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1NumericArrayConverter.html" target="_self">NumericArrayConverter</a></td><td class="desc">Conversion operator for Array </td></tr>
<tr id="row_0_78_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1NumericArrayConverter_3_01float_00_01half__t_00_012_00_01Round_01_4.html" target="_self">NumericArrayConverter&lt; float, half_t, 2, Round &gt;</a></td><td class="desc">Partial specialization for Array&lt;float, 2&gt; &lt;= Array&lt;half_t, 2&gt;, round to nearest </td></tr>
<tr id="row_0_79_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1NumericArrayConverter_3_01float_00_01half__t_00_01N_00_01Round_01_4.html" target="_self">NumericArrayConverter&lt; float, half_t, N, Round &gt;</a></td><td class="desc">Partial specialization for Array&lt;half&gt; &lt;= Array&lt;float&gt; </td></tr>
<tr id="row_0_80_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1NumericArrayConverter_3_01half__t_00_01float_00_012_00_01FloatRoundStyle_1_1round__to__nearest_01_4.html" target="_self">NumericArrayConverter&lt; half_t, float, 2, FloatRoundStyle::round_to_nearest &gt;</a></td><td class="desc">Partial specialization for Array&lt;half, 2&gt; &lt;= Array&lt;float, 2&gt;, round to nearest </td></tr>
<tr id="row_0_81_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1NumericArrayConverter_3_01half__t_00_01float_00_01N_00_01Round_01_4.html" target="_self">NumericArrayConverter&lt; half_t, float, N, Round &gt;</a></td><td class="desc">Partial specialization for Array&lt;half&gt; &lt;= Array&lt;float&gt; </td></tr>
<tr id="row_0_82_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1NumericConverter.html" target="_self">NumericConverter</a></td><td class="desc"></td></tr>
<tr id="row_0_83_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1NumericConverter_3_01float_00_01half__t_00_01Round_01_4.html" target="_self">NumericConverter&lt; float, half_t, Round &gt;</a></td><td class="desc">Partial specialization for float &lt;= <a class="el" href="structcutlass_1_1half__t.html" title="IEEE half-precision floating-point type. ">half_t</a> </td></tr>
<tr id="row_0_84_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1NumericConverter_3_01half__t_00_01float_00_01FloatRoundStyle_1_1round__to__nearest_01_4.html" target="_self">NumericConverter&lt; half_t, float, FloatRoundStyle::round_to_nearest &gt;</a></td><td class="desc">Specialization for round-to-nearest </td></tr>
<tr id="row_0_85_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1NumericConverter_3_01half__t_00_01float_00_01FloatRoundStyle_1_1round__toward__zero_01_4.html" target="_self">NumericConverter&lt; half_t, float, FloatRoundStyle::round_toward_zero &gt;</a></td><td class="desc">Specialization for round-toward-zero </td></tr>
<tr id="row_0_86_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1NumericConverter_3_01int8__t_00_01float_00_01Round_01_4.html" target="_self">NumericConverter&lt; int8_t, float, Round &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_87_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1NumericConverter_3_01T_00_01T_00_01Round_01_4.html" target="_self">NumericConverter&lt; T, T, Round &gt;</a></td><td class="desc">Partial specialization for float &lt;= <a class="el" href="structcutlass_1_1half__t.html" title="IEEE half-precision floating-point type. ">half_t</a> </td></tr>
<tr id="row_0_88_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1NumericConverterClamp.html" target="_self">NumericConverterClamp</a></td><td class="desc"></td></tr>
<tr id="row_0_89_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1plus.html" target="_self">plus</a></td><td class="desc"></td></tr>
<tr id="row_0_90_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1plus_3_01Array_3_01half__t_00_01N_01_4_01_4.html" target="_self">plus&lt; Array&lt; half_t, N &gt; &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_91_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1plus_3_01Array_3_01T_00_01N_01_4_01_4.html" target="_self">plus&lt; Array&lt; T, N &gt; &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_92_" style="display:none;"><td class="entry"><span style="width:16px;display:inline-block;">&#160;</span><span id="arr_0_92_" class="arrow" onclick="toggleFolder('0_92_')">&#9658;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1PredicateVector.html" target="_self">PredicateVector</a></td><td class="desc">Statically sized array of bits implementing </td></tr>
<tr id="row_0_92_0_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1PredicateVector_1_1ConstIterator.html" target="_self">ConstIterator</a></td><td class="desc">An iterator implementing <a class="el" href="group__predicate__iterator__concept.html">Predicate Iterator Concept</a> enabling sequential read and write access to predicates </td></tr>
<tr id="row_0_92_1_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1PredicateVector_1_1Iterator.html" target="_self">Iterator</a></td><td class="desc">An iterator implementing <a class="el" href="group__predicate__iterator__concept.html">Predicate Iterator Concept</a> enabling sequential read and write access to predicates </td></tr>
<tr id="row_0_92_2_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1PredicateVector_1_1TrivialIterator.html" target="_self">TrivialIterator</a></td><td class="desc"><a class="el" href="classcutlass_1_1PredicateVector_1_1Iterator.html" title="An iterator implementing Predicate Iterator Concept enabling sequential read and write access to pred...">Iterator</a> that always returns true </td></tr>
<tr id="row_0_93_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1RealType.html" target="_self">RealType</a></td><td class="desc">Used to determine the real-valued underlying type of a numeric type T </td></tr>
<tr id="row_0_94_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1RealType_3_01complex_3_01T_01_4_01_4.html" target="_self">RealType&lt; complex&lt; T &gt; &gt;</a></td><td class="desc">Partial specialization for complex-valued type </td></tr>
<tr id="row_0_95_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1ReferenceFactory.html" target="_self">ReferenceFactory</a></td><td class="desc"></td></tr>
<tr id="row_0_96_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1ReferenceFactory_3_01Element_00_01false_01_4.html" target="_self">ReferenceFactory&lt; Element, false &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_97_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1ReferenceFactory_3_01Element_00_01true_01_4.html" target="_self">ReferenceFactory&lt; Element, true &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_98_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1ScalarIO.html" target="_self">ScalarIO</a></td><td class="desc">Helper to enable formatted printing of CUTLASS scalar types to an ostream </td></tr>
<tr id="row_0_99_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1Semaphore.html" target="_self">Semaphore</a></td><td class="desc">CTA-wide semaphore for inter-CTA synchronization </td></tr>
<tr id="row_0_100_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1sizeof__bits.html" target="_self">sizeof_bits</a></td><td class="desc">Defines the size of an element in bits </td></tr>
<tr id="row_0_101_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1sizeof__bits_3_01Array_3_01T_00_01N_00_01RegisterSized_01_4_01_4.html" target="_self">sizeof_bits&lt; Array&lt; T, N, RegisterSized &gt; &gt;</a></td><td class="desc">Statically sized array for any data type </td></tr>
<tr id="row_0_102_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1sizeof__bits_3_01bin1__t_01_4.html" target="_self">sizeof_bits&lt; bin1_t &gt;</a></td><td class="desc">Defines the size of an element in bits - specialized for bin1_t </td></tr>
<tr id="row_0_103_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1sizeof__bits_3_01int4b__t_01_4.html" target="_self">sizeof_bits&lt; int4b_t &gt;</a></td><td class="desc">Defines the size of an element in bits - specialized for int4b_t </td></tr>
<tr id="row_0_104_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1sizeof__bits_3_01uint1b__t_01_4.html" target="_self">sizeof_bits&lt; uint1b_t &gt;</a></td><td class="desc">Defines the size of an element in bits - specialized for uint1b_t </td></tr>
<tr id="row_0_105_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1sizeof__bits_3_01uint4b__t_01_4.html" target="_self">sizeof_bits&lt; uint4b_t &gt;</a></td><td class="desc">Defines the size of an element in bits - specialized for uint4b_t </td></tr>
<tr id="row_0_106_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1sqrt__est.html" target="_self">sqrt_est</a></td><td class="desc"></td></tr>
<tr id="row_0_107_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1SubbyteReference.html" target="_self">SubbyteReference</a></td><td class="desc"></td></tr>
<tr id="row_0_108_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1Tensor4DCoord.html" target="_self">Tensor4DCoord</a></td><td class="desc">Defines a canonical 4D coordinate used by tensor operations </td></tr>
<tr id="row_0_109_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1TensorRef.html" target="_self">TensorRef</a></td><td class="desc"></td></tr>
<tr id="row_0_110_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="classcutlass_1_1TensorView.html" target="_self">TensorView</a></td><td class="desc"></td></tr>
<tr id="row_0_111_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1TypeTraits.html" target="_self">TypeTraits</a></td><td class="desc"></td></tr>
<tr id="row_0_112_" style="display:none;"><td class="entry"><span style="width:16px;display:inline-block;">&#160;</span><span id="arr_0_112_" class="arrow" onclick="toggleFolder('0_112_')">&#9658;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1TypeTraits_3_01complex_3_01double_01_4_01_4.html" target="_self">TypeTraits&lt; complex&lt; double &gt; &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_112_0_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1TypeTraits_3_01complex_3_01double_01_4_01_4_1_1integer__type.html" target="_self">integer_type</a></td><td class="desc"></td></tr>
<tr id="row_0_112_1_" style="display:none;"><td class="entry"><span style="width:48px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1TypeTraits_3_01complex_3_01double_01_4_01_4_1_1unsigned__type.html" target="_self">unsigned_type</a></td><td class="desc"></td></tr>
<tr id="row_0_113_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1TypeTraits_3_01complex_3_01float_01_4_01_4.html" target="_self">TypeTraits&lt; complex&lt; float &gt; &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_114_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1TypeTraits_3_01complex_3_01half_01_4_01_4.html" target="_self">TypeTraits&lt; complex&lt; half &gt; &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_115_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1TypeTraits_3_01complex_3_01half__t_01_4_01_4.html" target="_self">TypeTraits&lt; complex&lt; half_t &gt; &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_116_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1TypeTraits_3_01double_01_4.html" target="_self">TypeTraits&lt; double &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_117_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1TypeTraits_3_01float_01_4.html" target="_self">TypeTraits&lt; float &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_118_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1TypeTraits_3_01half__t_01_4.html" target="_self">TypeTraits&lt; half_t &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_119_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1TypeTraits_3_01int_01_4.html" target="_self">TypeTraits&lt; int &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_120_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1TypeTraits_3_01int64__t_01_4.html" target="_self">TypeTraits&lt; int64_t &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_121_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1TypeTraits_3_01int8__t_01_4.html" target="_self">TypeTraits&lt; int8_t &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_122_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1TypeTraits_3_01uint64__t_01_4.html" target="_self">TypeTraits&lt; uint64_t &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_123_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1TypeTraits_3_01uint8__t_01_4.html" target="_self">TypeTraits&lt; uint8_t &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_124_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1TypeTraits_3_01unsigned_01_4.html" target="_self">TypeTraits&lt; unsigned &gt;</a></td><td class="desc"></td></tr>
<tr id="row_0_125_" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structcutlass_1_1xor__add.html" target="_self">xor_add</a></td><td class="desc">Fused multiply-add </td></tr>
<tr id="row_1_"><td class="entry"><span style="width:0px;display:inline-block;">&#160;</span><span id="arr_1_" class="arrow" onclick="toggleFolder('1_')">&#9658;</span><span class="icona"><span class="icon">N</span></span><b>std</b></td><td class="desc">STL namespace </td></tr>
<tr id="row_1_0_" class="even" style="display:none;"><td class="entry"><span style="width:32px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structstd_1_1numeric__limits_3_01cutlass_1_1half__t_01_4.html" target="_self">numeric_limits&lt; cutlass::half_t &gt;</a></td><td class="desc">Numeric limits </td></tr>
<tr id="row_2_" class="even"><td class="entry"><span style="width:16px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structDebugType.html" target="_self">DebugType</a></td><td class="desc"></td></tr>
<tr id="row_3_"><td class="entry"><span style="width:16px;display:inline-block;">&#160;</span><span class="icona"><span class="icon">C</span></span><a class="el" href="structDebugValue.html" target="_self">DebugValue</a></td><td class="desc"></td></tr>
</table>
</div><!-- directory -->
</div><!-- contents -->
<!-- start footer part -->
<hr class="footer"/><address class="footer"><small>
Generated by &#160;<a href="http://www.doxygen.org/index.html">
<img class="footer" src="doxygen.png" alt="doxygen"/>
</a> 1.8.11
</small></address>
</body>
</html>

156
docs/arch_2mma_8h.html Normal file
View File

@ -0,0 +1,156 @@
<!DOCTYPE html PUBLIC "-//W3C//DTD XHTML 1.0 Transitional//EN" "http://www.w3.org/TR/xhtml1/DTD/xhtml1-transitional.dtd">
<html xmlns="http://www.w3.org/1999/xhtml">
<head>
<meta http-equiv="Content-Type" content="text/xhtml;charset=UTF-8"/>
<meta http-equiv="X-UA-Compatible" content="IE=9"/>
<meta name="generator" content="Doxygen 1.8.11"/>
<title>CUTLASS: mma.h File Reference</title>
<link href="tabs.css" rel="stylesheet" type="text/css"/>
<script type="text/javascript" src="jquery.js"></script>
<script type="text/javascript" src="dynsections.js"></script>
<link href="search/search.css" rel="stylesheet" type="text/css"/>
<script type="text/javascript" src="search/searchdata.js"></script>
<script type="text/javascript" src="search/search.js"></script>
<script type="text/javascript">
$(document).ready(function() { init_search(); });
</script>
<script type="text/x-mathjax-config">
MathJax.Hub.Config({
extensions: ["tex2jax.js"],
jax: ["input/TeX","output/HTML-CSS"],
});
</script><script type="text/javascript" src="http://cdn.mathjax.org/mathjax/latest/MathJax.js"></script>
<link href="doxygen.css" rel="stylesheet" type="text/css" />
</head>
<body>
<div id="top"><!-- do not remove this div, it is closed by doxygen! -->
<div id="titlearea">
<table cellspacing="0" cellpadding="0">
<tbody>
<tr style="height: 56px;">
<td id="projectlogo"><img alt="Logo" src="cutlass-logo-small.png"/></td>
<td id="projectalign" style="padding-left: 0.5em;">
<div id="projectname">CUTLASS
</div>
<div id="projectbrief">CUDA Templates for Linear Algebra Subroutines and Solvers</div>
</td>
</tr>
</tbody>
</table>
</div>
<!-- end header part -->
<!-- Generated by Doxygen 1.8.11 -->
<script type="text/javascript">
var searchBox = new SearchBox("searchBox", "search",false,'Search');
</script>
<div id="navrow1" class="tabs">
<ul class="tablist">
<li><a href="index.html"><span>Main&#160;Page</span></a></li>
<li><a href="modules.html"><span>Modules</span></a></li>
<li><a href="namespaces.html"><span>Namespaces</span></a></li>
<li><a href="annotated.html"><span>Classes</span></a></li>
<li class="current"><a href="files.html"><span>Files</span></a></li>
<li>
<div id="MSearchBox" class="MSearchBoxInactive">
<span class="left">
<img id="MSearchSelect" src="search/mag_sel.png"
onmouseover="return searchBox.OnSearchSelectShow()"
onmouseout="return searchBox.OnSearchSelectHide()"
alt=""/>
<input type="text" id="MSearchField" value="Search" accesskey="S"
onfocus="searchBox.OnSearchFieldFocus(true)"
onblur="searchBox.OnSearchFieldFocus(false)"
onkeyup="searchBox.OnSearchFieldChange(event)"/>
</span><span class="right">
<a id="MSearchClose" href="javascript:searchBox.CloseResultsWindow()"><img id="MSearchCloseImg" border="0" src="search/close.png" alt=""/></a>
</span>
</div>
</li>
</ul>
</div>
<div id="navrow2" class="tabs2">
<ul class="tablist">
<li><a href="files.html"><span>File&#160;List</span></a></li>
<li><a href="globals.html"><span>File&#160;Members</span></a></li>
</ul>
</div>
<!-- window showing the filter options -->
<div id="MSearchSelectWindow"
onmouseover="return searchBox.OnSearchSelectShow()"
onmouseout="return searchBox.OnSearchSelectHide()"
onkeydown="return searchBox.OnSearchSelectKey(event)">
</div>
<!-- iframe showing the search results (closed by default) -->
<div id="MSearchResultsWindow">
<iframe src="javascript:void(0)" frameborder="0"
name="MSearchResults" id="MSearchResults">
</iframe>
</div>
<div id="nav-path" class="navpath">
<ul>
<li class="navelem"><a class="el" href="dir_d44c64559bbebec7f509842c48db8b23.html">include</a></li><li class="navelem"><a class="el" href="dir_6baf2bb612a2f0daa69af3101ede80a1.html">cutlass</a></li><li class="navelem"><a class="el" href="dir_048c1df36ab9c2efbb0733edba6291c9.html">arch</a></li> </ul>
</div>
</div><!-- top -->
<div class="header">
<div class="summary">
<a href="#nested-classes">Classes</a> &#124;
<a href="#namespaces">Namespaces</a> </div>
<div class="headertitle">
<div class="title">arch/mma.h File Reference</div> </div>
</div><!--header-->
<div class="contents">
<p>Templates exposing architecture support for multiply-add operations.
<a href="#details">More...</a></p>
<div class="textblock"><code>#include &quot;<a class="el" href="array_8h_source.html">cutlass/array.h</a>&quot;</code><br />
<code>#include &quot;<a class="el" href="numeric__types_8h_source.html">cutlass/numeric_types.h</a>&quot;</code><br />
<code>#include &quot;<a class="el" href="include_2cutlass_2gemm_2gemm_8h_source.html">cutlass/gemm/gemm.h</a>&quot;</code><br />
<code>#include &quot;<a class="el" href="arch_2mma__sm50_8h_source.html">cutlass/arch/mma_sm50.h</a>&quot;</code><br />
<code>#include &quot;<a class="el" href="arch_2mma__sm60_8h_source.html">cutlass/arch/mma_sm60.h</a>&quot;</code><br />
<code>#include &quot;<a class="el" href="arch_2mma__sm61_8h_source.html">cutlass/arch/mma_sm61.h</a>&quot;</code><br />
<code>#include &quot;<a class="el" href="mma__sm70_8h_source.html">cutlass/arch/mma_sm70.h</a>&quot;</code><br />
<code>#include &quot;<a class="el" href="mma__sm75_8h_source.html">cutlass/arch/mma_sm75.h</a>&quot;</code><br />
</div><div class="textblock"><div class="dynheader">
Include dependency graph for arch/mma.h:</div>
<div class="dyncontent">
<div class="center"><img src="arch_2mma_8h__incl.png" border="0" usemap="#mma_8h" alt=""/></div>
<map name="mma_8h" id="mma_8h">
</map>
</div>
</div><div class="textblock"><div class="dynheader">
This graph shows which files directly or indirectly include this file:</div>
<div class="dyncontent">
<div class="center"><img src="arch_2mma_8h__dep__incl.png" border="0" usemap="#mma_8hdep" alt=""/></div>
<map name="mma_8hdep" id="mma_8hdep">
</map>
</div>
</div>
<p><a href="arch_2mma_8h_source.html">Go to the source code of this file.</a></p>
<table class="memberdecls">
<tr class="heading"><td colspan="2"><h2 class="groupheader"><a name="nested-classes"></a>
Classes</h2></td></tr>
<tr class="memitem:"><td class="memItemLeft" align="right" valign="top">struct &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="structcutlass_1_1arch_1_1Mma.html">cutlass::arch::Mma&lt; Shape_, kThreads_, ElementA, LayoutA, ElementB, LayoutB, ElementC, LayoutC, Operator &gt;</a></td></tr>
<tr class="memdesc:"><td class="mdescLeft">&#160;</td><td class="mdescRight">Matrix multiply-add operation. <a href="structcutlass_1_1arch_1_1Mma.html#details">More...</a><br /></td></tr>
<tr class="separator:"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:"><td class="memItemLeft" align="right" valign="top">struct &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="structcutlass_1_1arch_1_1Mma_3_01gemm_1_1GemmShape_3_011_00_011_00_011_01_4_00_011_00_01ElementAb6e65b2cf5ede7f41cb070a767158dee.html">cutlass::arch::Mma&lt; gemm::GemmShape&lt; 1, 1, 1 &gt;, 1, ElementA, LayoutA, ElementB, LayoutB, ElementC, LayoutC, Operator &gt;</a></td></tr>
<tr class="memdesc:"><td class="mdescLeft">&#160;</td><td class="mdescRight">Matrix multiply-add operation - specialized for 1x1x1x1 matrix multiply operation. <a href="structcutlass_1_1arch_1_1Mma_3_01gemm_1_1GemmShape_3_011_00_011_00_011_01_4_00_011_00_01ElementAb6e65b2cf5ede7f41cb070a767158dee.html#details">More...</a><br /></td></tr>
<tr class="separator:"><td class="memSeparator" colspan="2">&#160;</td></tr>
</table><table class="memberdecls">
<tr class="heading"><td colspan="2"><h2 class="groupheader"><a name="namespaces"></a>
Namespaces</h2></td></tr>
<tr class="memitem:namespacecutlass"><td class="memItemLeft" align="right" valign="top"> &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="namespacecutlass.html">cutlass</a></td></tr>
<tr class="separator:"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:namespacecutlass_1_1arch"><td class="memItemLeft" align="right" valign="top"> &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="namespacecutlass_1_1arch.html">cutlass::arch</a></td></tr>
<tr class="separator:"><td class="memSeparator" colspan="2">&#160;</td></tr>
</table>
</div><!-- contents -->
<!-- start footer part -->
<hr class="footer"/><address class="footer"><small>
Generated by &#160;<a href="http://www.doxygen.org/index.html">
<img class="footer" src="doxygen.png" alt="doxygen"/>
</a> 1.8.11
</small></address>
</body>
</html>

View File

@ -0,0 +1 @@
7d16b59e6ba0442b8a275a213d5da3a6

View File

@ -0,0 +1 @@
d1fff3f9d55a262110aa6a456caa91e0

File diff suppressed because one or more lines are too long

View File

@ -0,0 +1,176 @@
<!DOCTYPE html PUBLIC "-//W3C//DTD XHTML 1.0 Transitional//EN" "http://www.w3.org/TR/xhtml1/DTD/xhtml1-transitional.dtd">
<html xmlns="http://www.w3.org/1999/xhtml">
<head>
<meta http-equiv="Content-Type" content="text/xhtml;charset=UTF-8"/>
<meta http-equiv="X-UA-Compatible" content="IE=9"/>
<meta name="generator" content="Doxygen 1.8.11"/>
<title>CUTLASS: mma_sm50.h File Reference</title>
<link href="tabs.css" rel="stylesheet" type="text/css"/>
<script type="text/javascript" src="jquery.js"></script>
<script type="text/javascript" src="dynsections.js"></script>
<link href="search/search.css" rel="stylesheet" type="text/css"/>
<script type="text/javascript" src="search/searchdata.js"></script>
<script type="text/javascript" src="search/search.js"></script>
<script type="text/javascript">
$(document).ready(function() { init_search(); });
</script>
<script type="text/x-mathjax-config">
MathJax.Hub.Config({
extensions: ["tex2jax.js"],
jax: ["input/TeX","output/HTML-CSS"],
});
</script><script type="text/javascript" src="http://cdn.mathjax.org/mathjax/latest/MathJax.js"></script>
<link href="doxygen.css" rel="stylesheet" type="text/css" />
</head>
<body>
<div id="top"><!-- do not remove this div, it is closed by doxygen! -->
<div id="titlearea">
<table cellspacing="0" cellpadding="0">
<tbody>
<tr style="height: 56px;">
<td id="projectlogo"><img alt="Logo" src="cutlass-logo-small.png"/></td>
<td id="projectalign" style="padding-left: 0.5em;">
<div id="projectname">CUTLASS
</div>
<div id="projectbrief">CUDA Templates for Linear Algebra Subroutines and Solvers</div>
</td>
</tr>
</tbody>
</table>
</div>
<!-- end header part -->
<!-- Generated by Doxygen 1.8.11 -->
<script type="text/javascript">
var searchBox = new SearchBox("searchBox", "search",false,'Search');
</script>
<div id="navrow1" class="tabs">
<ul class="tablist">
<li><a href="index.html"><span>Main&#160;Page</span></a></li>
<li><a href="modules.html"><span>Modules</span></a></li>
<li><a href="namespaces.html"><span>Namespaces</span></a></li>
<li><a href="annotated.html"><span>Classes</span></a></li>
<li class="current"><a href="files.html"><span>Files</span></a></li>
<li>
<div id="MSearchBox" class="MSearchBoxInactive">
<span class="left">
<img id="MSearchSelect" src="search/mag_sel.png"
onmouseover="return searchBox.OnSearchSelectShow()"
onmouseout="return searchBox.OnSearchSelectHide()"
alt=""/>
<input type="text" id="MSearchField" value="Search" accesskey="S"
onfocus="searchBox.OnSearchFieldFocus(true)"
onblur="searchBox.OnSearchFieldFocus(false)"
onkeyup="searchBox.OnSearchFieldChange(event)"/>
</span><span class="right">
<a id="MSearchClose" href="javascript:searchBox.CloseResultsWindow()"><img id="MSearchCloseImg" border="0" src="search/close.png" alt=""/></a>
</span>
</div>
</li>
</ul>
</div>
<div id="navrow2" class="tabs2">
<ul class="tablist">
<li><a href="files.html"><span>File&#160;List</span></a></li>
<li><a href="globals.html"><span>File&#160;Members</span></a></li>
</ul>
</div>
<!-- window showing the filter options -->
<div id="MSearchSelectWindow"
onmouseover="return searchBox.OnSearchSelectShow()"
onmouseout="return searchBox.OnSearchSelectHide()"
onkeydown="return searchBox.OnSearchSelectKey(event)">
</div>
<!-- iframe showing the search results (closed by default) -->
<div id="MSearchResultsWindow">
<iframe src="javascript:void(0)" frameborder="0"
name="MSearchResults" id="MSearchResults">
</iframe>
</div>
<div id="nav-path" class="navpath">
<ul>
<li class="navelem"><a class="el" href="dir_d44c64559bbebec7f509842c48db8b23.html">include</a></li><li class="navelem"><a class="el" href="dir_6baf2bb612a2f0daa69af3101ede80a1.html">cutlass</a></li><li class="navelem"><a class="el" href="dir_048c1df36ab9c2efbb0733edba6291c9.html">arch</a></li> </ul>
</div>
</div><!-- top -->
<div class="header">
<div class="summary">
<a href="#nested-classes">Classes</a> &#124;
<a href="#namespaces">Namespaces</a> </div>
<div class="headertitle">
<div class="title">arch/mma_sm50.h File Reference</div> </div>
</div><!--header-->
<div class="contents">
<p>Matrix multiply.
<a href="#details">More...</a></p>
<div class="textblock"><code>#include &quot;<a class="el" href="arch_2mma_8h_source.html">cutlass/arch/mma.h</a>&quot;</code><br />
<code>#include &quot;<a class="el" href="complex_8h_source.html">cutlass/complex.h</a>&quot;</code><br />
<code>#include &quot;<a class="el" href="layout_2matrix_8h_source.html">cutlass/layout/matrix.h</a>&quot;</code><br />
<code>#include &quot;<a class="el" href="include_2cutlass_2gemm_2gemm_8h_source.html">cutlass/gemm/gemm.h</a>&quot;</code><br />
</div><div class="textblock"><div class="dynheader">
Include dependency graph for arch/mma_sm50.h:</div>
<div class="dyncontent">
<div class="center"><img src="arch_2mma__sm50_8h__incl.png" border="0" usemap="#mma__sm50_8h" alt=""/></div>
<map name="mma__sm50_8h" id="mma__sm50_8h">
</map>
</div>
</div><div class="textblock"><div class="dynheader">
This graph shows which files directly or indirectly include this file:</div>
<div class="dyncontent">
<div class="center"><img src="arch_2mma__sm50_8h__dep__incl.png" border="0" usemap="#mma__sm50_8hdep" alt=""/></div>
<map name="mma__sm50_8hdep" id="mma__sm50_8hdep">
</map>
</div>
</div>
<p><a href="arch_2mma__sm50_8h_source.html">Go to the source code of this file.</a></p>
<table class="memberdecls">
<tr class="heading"><td colspan="2"><h2 class="groupheader"><a name="nested-classes"></a>
Classes</h2></td></tr>
<tr class="memitem:"><td class="memItemLeft" align="right" valign="top">struct &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="structcutlass_1_1arch_1_1Mma_3_01gemm_1_1GemmShape_3_011_00_011_00_011_01_4_00_011_00_01float_004bb3fd76ca2af7b3210676fa9644d95b.html">cutlass::arch::Mma&lt; gemm::GemmShape&lt; 1, 1, 1 &gt;, 1, float, LayoutA, float, LayoutB, float, LayoutC, OpMultiplyAdd &gt;</a></td></tr>
<tr class="memdesc:"><td class="mdescLeft">&#160;</td><td class="mdescRight">Matrix multiply-add operation. <a href="structcutlass_1_1arch_1_1Mma_3_01gemm_1_1GemmShape_3_011_00_011_00_011_01_4_00_011_00_01float_004bb3fd76ca2af7b3210676fa9644d95b.html#details">More...</a><br /></td></tr>
<tr class="separator:"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:"><td class="memItemLeft" align="right" valign="top">struct &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="structcutlass_1_1arch_1_1Mma_3_01gemm_1_1GemmShape_3_011_00_011_00_011_01_4_00_011_00_01double_0aa57e6a2e6b5da37d10688bf99419a23.html">cutlass::arch::Mma&lt; gemm::GemmShape&lt; 1, 1, 1 &gt;, 1, double, LayoutA, double, LayoutB, double, LayoutC, OpMultiplyAdd &gt;</a></td></tr>
<tr class="memdesc:"><td class="mdescLeft">&#160;</td><td class="mdescRight">Matrix multiply-add operation. <a href="structcutlass_1_1arch_1_1Mma_3_01gemm_1_1GemmShape_3_011_00_011_00_011_01_4_00_011_00_01double_0aa57e6a2e6b5da37d10688bf99419a23.html#details">More...</a><br /></td></tr>
<tr class="separator:"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:"><td class="memItemLeft" align="right" valign="top">struct &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="structcutlass_1_1arch_1_1Mma_3_01gemm_1_1GemmShape_3_011_00_011_00_011_01_4_00_011_00_01int_00_00b2dff9ce8caad9aff5bc6a355539161.html">cutlass::arch::Mma&lt; gemm::GemmShape&lt; 1, 1, 1 &gt;, 1, int, LayoutA, int, LayoutB, int, LayoutC, OpMultiplyAdd &gt;</a></td></tr>
<tr class="memdesc:"><td class="mdescLeft">&#160;</td><td class="mdescRight">Matrix multiply-add operation. <a href="structcutlass_1_1arch_1_1Mma_3_01gemm_1_1GemmShape_3_011_00_011_00_011_01_4_00_011_00_01int_00_00b2dff9ce8caad9aff5bc6a355539161.html#details">More...</a><br /></td></tr>
<tr class="separator:"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:"><td class="memItemLeft" align="right" valign="top">struct &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="structcutlass_1_1arch_1_1Mma_3_01gemm_1_1GemmShape_3_011_00_011_00_011_01_4_00_011_00_01complex_76f9d24016e1b4167b16f4d7628c9546.html">cutlass::arch::Mma&lt; gemm::GemmShape&lt; 1, 1, 1 &gt;, 1, complex&lt; float &gt;, LayoutA, complex&lt; float &gt;, LayoutB, complex&lt; float &gt;, LayoutC, OpMultiplyAdd &gt;</a></td></tr>
<tr class="memdesc:"><td class="mdescLeft">&#160;</td><td class="mdescRight">Matrix multiply-add operation. <a href="structcutlass_1_1arch_1_1Mma_3_01gemm_1_1GemmShape_3_011_00_011_00_011_01_4_00_011_00_01complex_76f9d24016e1b4167b16f4d7628c9546.html#details">More...</a><br /></td></tr>
<tr class="separator:"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:"><td class="memItemLeft" align="right" valign="top">struct &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="structcutlass_1_1arch_1_1Mma_3_01gemm_1_1GemmShape_3_011_00_011_00_011_01_4_00_011_00_01complex_f1c9d2ee842455cd0c5b71d56108d468.html">cutlass::arch::Mma&lt; gemm::GemmShape&lt; 1, 1, 1 &gt;, 1, complex&lt; float &gt;, LayoutA, float, LayoutB, complex&lt; float &gt;, LayoutC, OpMultiplyAdd &gt;</a></td></tr>
<tr class="memdesc:"><td class="mdescLeft">&#160;</td><td class="mdescRight">Matrix multiply-add operation. <a href="structcutlass_1_1arch_1_1Mma_3_01gemm_1_1GemmShape_3_011_00_011_00_011_01_4_00_011_00_01complex_f1c9d2ee842455cd0c5b71d56108d468.html#details">More...</a><br /></td></tr>
<tr class="separator:"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:"><td class="memItemLeft" align="right" valign="top">struct &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="structcutlass_1_1arch_1_1Mma_3_01gemm_1_1GemmShape_3_011_00_011_00_011_01_4_00_011_00_01float_00e3e12e263df6506b8cf06c3f4d478b8e.html">cutlass::arch::Mma&lt; gemm::GemmShape&lt; 1, 1, 1 &gt;, 1, float, LayoutA, complex&lt; float &gt;, LayoutB, complex&lt; float &gt;, LayoutC, OpMultiplyAdd &gt;</a></td></tr>
<tr class="memdesc:"><td class="mdescLeft">&#160;</td><td class="mdescRight">Matrix multiply-add operation. <a href="structcutlass_1_1arch_1_1Mma_3_01gemm_1_1GemmShape_3_011_00_011_00_011_01_4_00_011_00_01float_00e3e12e263df6506b8cf06c3f4d478b8e.html#details">More...</a><br /></td></tr>
<tr class="separator:"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:"><td class="memItemLeft" align="right" valign="top">struct &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="structcutlass_1_1arch_1_1Mma_3_01gemm_1_1GemmShape_3_011_00_011_00_011_01_4_00_011_00_01complex_30fa42e1ad201df010637cd22fc070a1.html">cutlass::arch::Mma&lt; gemm::GemmShape&lt; 1, 1, 1 &gt;, 1, complex&lt; double &gt;, LayoutA, complex&lt; double &gt;, LayoutB, complex&lt; double &gt;, LayoutC, OpMultiplyAdd &gt;</a></td></tr>
<tr class="memdesc:"><td class="mdescLeft">&#160;</td><td class="mdescRight">Matrix multiply-add operation. <a href="structcutlass_1_1arch_1_1Mma_3_01gemm_1_1GemmShape_3_011_00_011_00_011_01_4_00_011_00_01complex_30fa42e1ad201df010637cd22fc070a1.html#details">More...</a><br /></td></tr>
<tr class="separator:"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:"><td class="memItemLeft" align="right" valign="top">struct &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="structcutlass_1_1arch_1_1Mma_3_01gemm_1_1GemmShape_3_011_00_011_00_011_01_4_00_011_00_01complex_48b3a43bc03fff93a111ac01abe7e40d.html">cutlass::arch::Mma&lt; gemm::GemmShape&lt; 1, 1, 1 &gt;, 1, complex&lt; double &gt;, LayoutA, double, LayoutB, complex&lt; double &gt;, LayoutC, OpMultiplyAdd &gt;</a></td></tr>
<tr class="memdesc:"><td class="mdescLeft">&#160;</td><td class="mdescRight">Matrix multiply-add operation. <a href="structcutlass_1_1arch_1_1Mma_3_01gemm_1_1GemmShape_3_011_00_011_00_011_01_4_00_011_00_01complex_48b3a43bc03fff93a111ac01abe7e40d.html#details">More...</a><br /></td></tr>
<tr class="separator:"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:"><td class="memItemLeft" align="right" valign="top">struct &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="structcutlass_1_1arch_1_1Mma_3_01gemm_1_1GemmShape_3_011_00_011_00_011_01_4_00_011_00_01double_070b94670e040ed5855e5b42d5ca8a443.html">cutlass::arch::Mma&lt; gemm::GemmShape&lt; 1, 1, 1 &gt;, 1, double, LayoutA, complex&lt; double &gt;, LayoutB, complex&lt; double &gt;, LayoutC, OpMultiplyAdd &gt;</a></td></tr>
<tr class="memdesc:"><td class="mdescLeft">&#160;</td><td class="mdescRight">Matrix multiply-add operation. <a href="structcutlass_1_1arch_1_1Mma_3_01gemm_1_1GemmShape_3_011_00_011_00_011_01_4_00_011_00_01double_070b94670e040ed5855e5b42d5ca8a443.html#details">More...</a><br /></td></tr>
<tr class="separator:"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:"><td class="memItemLeft" align="right" valign="top">struct &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="structcutlass_1_1arch_1_1Mma_3_01gemm_1_1GemmShape_3_011_00_011_00_011_01_4_00_011_00_01half__t_4f30ee91f7bb3844ff7579c68d078818.html">cutlass::arch::Mma&lt; gemm::GemmShape&lt; 1, 1, 1 &gt;, 1, half_t, LayoutA, half_t, LayoutB, float, LayoutC, OpMultiplyAdd &gt;</a></td></tr>
<tr class="memdesc:"><td class="mdescLeft">&#160;</td><td class="mdescRight">Matrix multiply-add operation. <a href="structcutlass_1_1arch_1_1Mma_3_01gemm_1_1GemmShape_3_011_00_011_00_011_01_4_00_011_00_01half__t_4f30ee91f7bb3844ff7579c68d078818.html#details">More...</a><br /></td></tr>
<tr class="separator:"><td class="memSeparator" colspan="2">&#160;</td></tr>
</table><table class="memberdecls">
<tr class="heading"><td colspan="2"><h2 class="groupheader"><a name="namespaces"></a>
Namespaces</h2></td></tr>
<tr class="memitem:namespacecutlass"><td class="memItemLeft" align="right" valign="top"> &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="namespacecutlass.html">cutlass</a></td></tr>
<tr class="separator:"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:namespacecutlass_1_1arch"><td class="memItemLeft" align="right" valign="top"> &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="namespacecutlass_1_1arch.html">cutlass::arch</a></td></tr>
<tr class="separator:"><td class="memSeparator" colspan="2">&#160;</td></tr>
</table>
</div><!-- contents -->
<!-- start footer part -->
<hr class="footer"/><address class="footer"><small>
Generated by &#160;<a href="http://www.doxygen.org/index.html">
<img class="footer" src="doxygen.png" alt="doxygen"/>
</a> 1.8.11
</small></address>
</body>
</html>

View File

@ -0,0 +1 @@
988e6466c703c4e63c9a889b8c3c54b5

View File

@ -0,0 +1 @@
03f1613fdffbd6e7575de0d2967d08bf

File diff suppressed because one or more lines are too long

View File

@ -0,0 +1,157 @@
<!DOCTYPE html PUBLIC "-//W3C//DTD XHTML 1.0 Transitional//EN" "http://www.w3.org/TR/xhtml1/DTD/xhtml1-transitional.dtd">
<html xmlns="http://www.w3.org/1999/xhtml">
<head>
<meta http-equiv="Content-Type" content="text/xhtml;charset=UTF-8"/>
<meta http-equiv="X-UA-Compatible" content="IE=9"/>
<meta name="generator" content="Doxygen 1.8.11"/>
<title>CUTLASS: mma_sm60.h File Reference</title>
<link href="tabs.css" rel="stylesheet" type="text/css"/>
<script type="text/javascript" src="jquery.js"></script>
<script type="text/javascript" src="dynsections.js"></script>
<link href="search/search.css" rel="stylesheet" type="text/css"/>
<script type="text/javascript" src="search/searchdata.js"></script>
<script type="text/javascript" src="search/search.js"></script>
<script type="text/javascript">
$(document).ready(function() { init_search(); });
</script>
<script type="text/x-mathjax-config">
MathJax.Hub.Config({
extensions: ["tex2jax.js"],
jax: ["input/TeX","output/HTML-CSS"],
});
</script><script type="text/javascript" src="http://cdn.mathjax.org/mathjax/latest/MathJax.js"></script>
<link href="doxygen.css" rel="stylesheet" type="text/css" />
</head>
<body>
<div id="top"><!-- do not remove this div, it is closed by doxygen! -->
<div id="titlearea">
<table cellspacing="0" cellpadding="0">
<tbody>
<tr style="height: 56px;">
<td id="projectlogo"><img alt="Logo" src="cutlass-logo-small.png"/></td>
<td id="projectalign" style="padding-left: 0.5em;">
<div id="projectname">CUTLASS
</div>
<div id="projectbrief">CUDA Templates for Linear Algebra Subroutines and Solvers</div>
</td>
</tr>
</tbody>
</table>
</div>
<!-- end header part -->
<!-- Generated by Doxygen 1.8.11 -->
<script type="text/javascript">
var searchBox = new SearchBox("searchBox", "search",false,'Search');
</script>
<div id="navrow1" class="tabs">
<ul class="tablist">
<li><a href="index.html"><span>Main&#160;Page</span></a></li>
<li><a href="modules.html"><span>Modules</span></a></li>
<li><a href="namespaces.html"><span>Namespaces</span></a></li>
<li><a href="annotated.html"><span>Classes</span></a></li>
<li class="current"><a href="files.html"><span>Files</span></a></li>
<li>
<div id="MSearchBox" class="MSearchBoxInactive">
<span class="left">
<img id="MSearchSelect" src="search/mag_sel.png"
onmouseover="return searchBox.OnSearchSelectShow()"
onmouseout="return searchBox.OnSearchSelectHide()"
alt=""/>
<input type="text" id="MSearchField" value="Search" accesskey="S"
onfocus="searchBox.OnSearchFieldFocus(true)"
onblur="searchBox.OnSearchFieldFocus(false)"
onkeyup="searchBox.OnSearchFieldChange(event)"/>
</span><span class="right">
<a id="MSearchClose" href="javascript:searchBox.CloseResultsWindow()"><img id="MSearchCloseImg" border="0" src="search/close.png" alt=""/></a>
</span>
</div>
</li>
</ul>
</div>
<div id="navrow2" class="tabs2">
<ul class="tablist">
<li><a href="files.html"><span>File&#160;List</span></a></li>
<li><a href="globals.html"><span>File&#160;Members</span></a></li>
</ul>
</div>
<!-- window showing the filter options -->
<div id="MSearchSelectWindow"
onmouseover="return searchBox.OnSearchSelectShow()"
onmouseout="return searchBox.OnSearchSelectHide()"
onkeydown="return searchBox.OnSearchSelectKey(event)">
</div>
<!-- iframe showing the search results (closed by default) -->
<div id="MSearchResultsWindow">
<iframe src="javascript:void(0)" frameborder="0"
name="MSearchResults" id="MSearchResults">
</iframe>
</div>
<div id="nav-path" class="navpath">
<ul>
<li class="navelem"><a class="el" href="dir_d44c64559bbebec7f509842c48db8b23.html">include</a></li><li class="navelem"><a class="el" href="dir_6baf2bb612a2f0daa69af3101ede80a1.html">cutlass</a></li><li class="navelem"><a class="el" href="dir_048c1df36ab9c2efbb0733edba6291c9.html">arch</a></li> </ul>
</div>
</div><!-- top -->
<div class="header">
<div class="summary">
<a href="#nested-classes">Classes</a> &#124;
<a href="#namespaces">Namespaces</a> </div>
<div class="headertitle">
<div class="title">arch/mma_sm60.h File Reference</div> </div>
</div><!--header-->
<div class="contents">
<p>Matrix multiply.
<a href="#details">More...</a></p>
<div class="textblock"><code>#include &lt;cuda_fp16.h&gt;</code><br />
<code>#include &quot;<a class="el" href="arch_2mma_8h_source.html">cutlass/arch/mma.h</a>&quot;</code><br />
<code>#include &quot;<a class="el" href="layout_2matrix_8h_source.html">cutlass/layout/matrix.h</a>&quot;</code><br />
</div><div class="textblock"><div class="dynheader">
Include dependency graph for arch/mma_sm60.h:</div>
<div class="dyncontent">
<div class="center"><img src="arch_2mma__sm60_8h__incl.png" border="0" usemap="#mma__sm60_8h" alt=""/></div>
<map name="mma__sm60_8h" id="mma__sm60_8h">
</map>
</div>
</div><div class="textblock"><div class="dynheader">
This graph shows which files directly or indirectly include this file:</div>
<div class="dyncontent">
<div class="center"><img src="arch_2mma__sm60_8h__dep__incl.png" border="0" usemap="#mma__sm60_8hdep" alt=""/></div>
<map name="mma__sm60_8hdep" id="mma__sm60_8hdep">
</map>
</div>
</div>
<p><a href="arch_2mma__sm60_8h_source.html">Go to the source code of this file.</a></p>
<table class="memberdecls">
<tr class="heading"><td colspan="2"><h2 class="groupheader"><a name="nested-classes"></a>
Classes</h2></td></tr>
<tr class="memitem:"><td class="memItemLeft" align="right" valign="top">struct &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="structcutlass_1_1arch_1_1Mma_3_01gemm_1_1GemmShape_3_012_00_011_00_011_01_4_00_011_00_01half__t_8cf78649807b93684f3d431bfa34ee28.html">cutlass::arch::Mma&lt; gemm::GemmShape&lt; 2, 1, 1 &gt;, 1, half_t, LayoutA, half_t, LayoutB, half_t, LayoutC, OpMultiplyAdd &gt;</a></td></tr>
<tr class="memdesc:"><td class="mdescLeft">&#160;</td><td class="mdescRight">Matrix multiply-add operation. <a href="structcutlass_1_1arch_1_1Mma_3_01gemm_1_1GemmShape_3_012_00_011_00_011_01_4_00_011_00_01half__t_8cf78649807b93684f3d431bfa34ee28.html#details">More...</a><br /></td></tr>
<tr class="separator:"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:"><td class="memItemLeft" align="right" valign="top">struct &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="structcutlass_1_1arch_1_1Mma_3_01gemm_1_1GemmShape_3_011_00_012_00_011_01_4_00_011_00_01half__t_f3dc2e59f857ada163d1e0781ea8f391.html">cutlass::arch::Mma&lt; gemm::GemmShape&lt; 1, 2, 1 &gt;, 1, half_t, LayoutA, half_t, LayoutB, half_t, layout::RowMajor, OpMultiplyAdd &gt;</a></td></tr>
<tr class="memdesc:"><td class="mdescLeft">&#160;</td><td class="mdescRight">Matrix multiply-add operation. <a href="structcutlass_1_1arch_1_1Mma_3_01gemm_1_1GemmShape_3_011_00_012_00_011_01_4_00_011_00_01half__t_f3dc2e59f857ada163d1e0781ea8f391.html#details">More...</a><br /></td></tr>
<tr class="separator:"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:"><td class="memItemLeft" align="right" valign="top">struct &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="structcutlass_1_1arch_1_1Mma_3_01gemm_1_1GemmShape_3_012_00_012_00_011_01_4_00_011_00_01half__t_ccde11d1bbbdab3702772ce44eb9729a.html">cutlass::arch::Mma&lt; gemm::GemmShape&lt; 2, 2, 1 &gt;, 1, half_t, layout::ColumnMajor, half_t, layout::RowMajor, half_t, layout::ColumnMajor, OpMultiplyAdd &gt;</a></td></tr>
<tr class="memdesc:"><td class="mdescLeft">&#160;</td><td class="mdescRight">Matrix multiply-add operation. <a href="structcutlass_1_1arch_1_1Mma_3_01gemm_1_1GemmShape_3_012_00_012_00_011_01_4_00_011_00_01half__t_ccde11d1bbbdab3702772ce44eb9729a.html#details">More...</a><br /></td></tr>
<tr class="separator:"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:"><td class="memItemLeft" align="right" valign="top">struct &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="structcutlass_1_1arch_1_1Mma_3_01gemm_1_1GemmShape_3_012_00_012_00_011_01_4_00_011_00_01half__t_c07cc6439298fa5486a719e577be2538.html">cutlass::arch::Mma&lt; gemm::GemmShape&lt; 2, 2, 1 &gt;, 1, half_t, layout::ColumnMajor, half_t, layout::RowMajor, half_t, layout::RowMajor, OpMultiplyAdd &gt;</a></td></tr>
<tr class="memdesc:"><td class="mdescLeft">&#160;</td><td class="mdescRight">Matrix multiply-add operation. <a href="structcutlass_1_1arch_1_1Mma_3_01gemm_1_1GemmShape_3_012_00_012_00_011_01_4_00_011_00_01half__t_c07cc6439298fa5486a719e577be2538.html#details">More...</a><br /></td></tr>
<tr class="separator:"><td class="memSeparator" colspan="2">&#160;</td></tr>
</table><table class="memberdecls">
<tr class="heading"><td colspan="2"><h2 class="groupheader"><a name="namespaces"></a>
Namespaces</h2></td></tr>
<tr class="memitem:namespacecutlass"><td class="memItemLeft" align="right" valign="top"> &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="namespacecutlass.html">cutlass</a></td></tr>
<tr class="separator:"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:namespacecutlass_1_1arch"><td class="memItemLeft" align="right" valign="top"> &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="namespacecutlass_1_1arch.html">cutlass::arch</a></td></tr>
<tr class="separator:"><td class="memSeparator" colspan="2">&#160;</td></tr>
</table>
</div><!-- contents -->
<!-- start footer part -->
<hr class="footer"/><address class="footer"><small>
Generated by &#160;<a href="http://www.doxygen.org/index.html">
<img class="footer" src="doxygen.png" alt="doxygen"/>
</a> 1.8.11
</small></address>
</body>
</html>

View File

@ -0,0 +1 @@
ba69b14e3936946092854211499ae9fa

View File

@ -0,0 +1 @@
e820099c55f2397639bb210d76ec4c05

File diff suppressed because one or more lines are too long

View File

@ -0,0 +1,149 @@
<!DOCTYPE html PUBLIC "-//W3C//DTD XHTML 1.0 Transitional//EN" "http://www.w3.org/TR/xhtml1/DTD/xhtml1-transitional.dtd">
<html xmlns="http://www.w3.org/1999/xhtml">
<head>
<meta http-equiv="Content-Type" content="text/xhtml;charset=UTF-8"/>
<meta http-equiv="X-UA-Compatible" content="IE=9"/>
<meta name="generator" content="Doxygen 1.8.11"/>
<title>CUTLASS: mma_sm61.h File Reference</title>
<link href="tabs.css" rel="stylesheet" type="text/css"/>
<script type="text/javascript" src="jquery.js"></script>
<script type="text/javascript" src="dynsections.js"></script>
<link href="search/search.css" rel="stylesheet" type="text/css"/>
<script type="text/javascript" src="search/searchdata.js"></script>
<script type="text/javascript" src="search/search.js"></script>
<script type="text/javascript">
$(document).ready(function() { init_search(); });
</script>
<script type="text/x-mathjax-config">
MathJax.Hub.Config({
extensions: ["tex2jax.js"],
jax: ["input/TeX","output/HTML-CSS"],
});
</script><script type="text/javascript" src="http://cdn.mathjax.org/mathjax/latest/MathJax.js"></script>
<link href="doxygen.css" rel="stylesheet" type="text/css" />
</head>
<body>
<div id="top"><!-- do not remove this div, it is closed by doxygen! -->
<div id="titlearea">
<table cellspacing="0" cellpadding="0">
<tbody>
<tr style="height: 56px;">
<td id="projectlogo"><img alt="Logo" src="cutlass-logo-small.png"/></td>
<td id="projectalign" style="padding-left: 0.5em;">
<div id="projectname">CUTLASS
</div>
<div id="projectbrief">CUDA Templates for Linear Algebra Subroutines and Solvers</div>
</td>
</tr>
</tbody>
</table>
</div>
<!-- end header part -->
<!-- Generated by Doxygen 1.8.11 -->
<script type="text/javascript">
var searchBox = new SearchBox("searchBox", "search",false,'Search');
</script>
<div id="navrow1" class="tabs">
<ul class="tablist">
<li><a href="index.html"><span>Main&#160;Page</span></a></li>
<li><a href="modules.html"><span>Modules</span></a></li>
<li><a href="namespaces.html"><span>Namespaces</span></a></li>
<li><a href="annotated.html"><span>Classes</span></a></li>
<li class="current"><a href="files.html"><span>Files</span></a></li>
<li>
<div id="MSearchBox" class="MSearchBoxInactive">
<span class="left">
<img id="MSearchSelect" src="search/mag_sel.png"
onmouseover="return searchBox.OnSearchSelectShow()"
onmouseout="return searchBox.OnSearchSelectHide()"
alt=""/>
<input type="text" id="MSearchField" value="Search" accesskey="S"
onfocus="searchBox.OnSearchFieldFocus(true)"
onblur="searchBox.OnSearchFieldFocus(false)"
onkeyup="searchBox.OnSearchFieldChange(event)"/>
</span><span class="right">
<a id="MSearchClose" href="javascript:searchBox.CloseResultsWindow()"><img id="MSearchCloseImg" border="0" src="search/close.png" alt=""/></a>
</span>
</div>
</li>
</ul>
</div>
<div id="navrow2" class="tabs2">
<ul class="tablist">
<li><a href="files.html"><span>File&#160;List</span></a></li>
<li><a href="globals.html"><span>File&#160;Members</span></a></li>
</ul>
</div>
<!-- window showing the filter options -->
<div id="MSearchSelectWindow"
onmouseover="return searchBox.OnSearchSelectShow()"
onmouseout="return searchBox.OnSearchSelectHide()"
onkeydown="return searchBox.OnSearchSelectKey(event)">
</div>
<!-- iframe showing the search results (closed by default) -->
<div id="MSearchResultsWindow">
<iframe src="javascript:void(0)" frameborder="0"
name="MSearchResults" id="MSearchResults">
</iframe>
</div>
<div id="nav-path" class="navpath">
<ul>
<li class="navelem"><a class="el" href="dir_d44c64559bbebec7f509842c48db8b23.html">include</a></li><li class="navelem"><a class="el" href="dir_6baf2bb612a2f0daa69af3101ede80a1.html">cutlass</a></li><li class="navelem"><a class="el" href="dir_048c1df36ab9c2efbb0733edba6291c9.html">arch</a></li> </ul>
</div>
</div><!-- top -->
<div class="header">
<div class="summary">
<a href="#nested-classes">Classes</a> &#124;
<a href="#namespaces">Namespaces</a> </div>
<div class="headertitle">
<div class="title">arch/mma_sm61.h File Reference</div> </div>
</div><!--header-->
<div class="contents">
<p>Matrix multiply.
<a href="#details">More...</a></p>
<div class="textblock"><code>#include &quot;<a class="el" href="layout_2matrix_8h_source.html">cutlass/layout/matrix.h</a>&quot;</code><br />
</div><div class="textblock"><div class="dynheader">
Include dependency graph for arch/mma_sm61.h:</div>
<div class="dyncontent">
<div class="center"><img src="arch_2mma__sm61_8h__incl.png" border="0" usemap="#mma__sm61_8h" alt=""/></div>
<map name="mma__sm61_8h" id="mma__sm61_8h">
</map>
</div>
</div><div class="textblock"><div class="dynheader">
This graph shows which files directly or indirectly include this file:</div>
<div class="dyncontent">
<div class="center"><img src="arch_2mma__sm61_8h__dep__incl.png" border="0" usemap="#mma__sm61_8hdep" alt=""/></div>
<map name="mma__sm61_8hdep" id="mma__sm61_8hdep">
</map>
</div>
</div>
<p><a href="arch_2mma__sm61_8h_source.html">Go to the source code of this file.</a></p>
<table class="memberdecls">
<tr class="heading"><td colspan="2"><h2 class="groupheader"><a name="nested-classes"></a>
Classes</h2></td></tr>
<tr class="memitem:"><td class="memItemLeft" align="right" valign="top">struct &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="structcutlass_1_1arch_1_1Mma_3_01gemm_1_1GemmShape_3_011_00_011_00_014_01_4_00_011_00_01int8__t_a1ef6624fc8c10126f17f4ee88283d72.html">cutlass::arch::Mma&lt; gemm::GemmShape&lt; 1, 1, 4 &gt;, 1, int8_t, LayoutA, int8_t, LayoutB, int, LayoutC, OpMultiplyAdd &gt;</a></td></tr>
<tr class="memdesc:"><td class="mdescLeft">&#160;</td><td class="mdescRight">Matrix multiply-add operation. <a href="structcutlass_1_1arch_1_1Mma_3_01gemm_1_1GemmShape_3_011_00_011_00_014_01_4_00_011_00_01int8__t_a1ef6624fc8c10126f17f4ee88283d72.html#details">More...</a><br /></td></tr>
<tr class="separator:"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:"><td class="memItemLeft" align="right" valign="top">struct &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="structcutlass_1_1arch_1_1Mma_3_01gemm_1_1GemmShape_3_011_00_011_00_012_01_4_00_011_00_01int16__t8c4bac365710598317a69c489f7239db.html">cutlass::arch::Mma&lt; gemm::GemmShape&lt; 1, 1, 2 &gt;, 1, int16_t, layout::RowMajor, int16_t, layout::ColumnMajor, int, LayoutC, OpMultiplyAdd &gt;</a></td></tr>
<tr class="memdesc:"><td class="mdescLeft">&#160;</td><td class="mdescRight">Matrix multiply-add operation. <a href="structcutlass_1_1arch_1_1Mma_3_01gemm_1_1GemmShape_3_011_00_011_00_012_01_4_00_011_00_01int16__t8c4bac365710598317a69c489f7239db.html#details">More...</a><br /></td></tr>
<tr class="separator:"><td class="memSeparator" colspan="2">&#160;</td></tr>
</table><table class="memberdecls">
<tr class="heading"><td colspan="2"><h2 class="groupheader"><a name="namespaces"></a>
Namespaces</h2></td></tr>
<tr class="memitem:namespacecutlass"><td class="memItemLeft" align="right" valign="top"> &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="namespacecutlass.html">cutlass</a></td></tr>
<tr class="separator:"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:namespacecutlass_1_1arch"><td class="memItemLeft" align="right" valign="top"> &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="namespacecutlass_1_1arch.html">cutlass::arch</a></td></tr>
<tr class="separator:"><td class="memSeparator" colspan="2">&#160;</td></tr>
</table>
</div><!-- contents -->
<!-- start footer part -->
<hr class="footer"/><address class="footer"><small>
Generated by &#160;<a href="http://www.doxygen.org/index.html">
<img class="footer" src="doxygen.png" alt="doxygen"/>
</a> 1.8.11
</small></address>
</body>
</html>

View File

@ -0,0 +1 @@
1faaf1631d5f0e44d6cc6c7121e6972e

View File

@ -0,0 +1 @@
8cce8aef2d98c4082d68734b538253c7

File diff suppressed because one or more lines are too long

147
docs/arch_8h.html Normal file
View File

@ -0,0 +1,147 @@
<!DOCTYPE html PUBLIC "-//W3C//DTD XHTML 1.0 Transitional//EN" "http://www.w3.org/TR/xhtml1/DTD/xhtml1-transitional.dtd">
<html xmlns="http://www.w3.org/1999/xhtml">
<head>
<meta http-equiv="Content-Type" content="text/xhtml;charset=UTF-8"/>
<meta http-equiv="X-UA-Compatible" content="IE=9"/>
<meta name="generator" content="Doxygen 1.8.11"/>
<title>CUTLASS: arch.h File Reference</title>
<link href="tabs.css" rel="stylesheet" type="text/css"/>
<script type="text/javascript" src="jquery.js"></script>
<script type="text/javascript" src="dynsections.js"></script>
<link href="search/search.css" rel="stylesheet" type="text/css"/>
<script type="text/javascript" src="search/searchdata.js"></script>
<script type="text/javascript" src="search/search.js"></script>
<script type="text/javascript">
$(document).ready(function() { init_search(); });
</script>
<script type="text/x-mathjax-config">
MathJax.Hub.Config({
extensions: ["tex2jax.js"],
jax: ["input/TeX","output/HTML-CSS"],
});
</script><script type="text/javascript" src="http://cdn.mathjax.org/mathjax/latest/MathJax.js"></script>
<link href="doxygen.css" rel="stylesheet" type="text/css" />
</head>
<body>
<div id="top"><!-- do not remove this div, it is closed by doxygen! -->
<div id="titlearea">
<table cellspacing="0" cellpadding="0">
<tbody>
<tr style="height: 56px;">
<td id="projectlogo"><img alt="Logo" src="cutlass-logo-small.png"/></td>
<td id="projectalign" style="padding-left: 0.5em;">
<div id="projectname">CUTLASS
</div>
<div id="projectbrief">CUDA Templates for Linear Algebra Subroutines and Solvers</div>
</td>
</tr>
</tbody>
</table>
</div>
<!-- end header part -->
<!-- Generated by Doxygen 1.8.11 -->
<script type="text/javascript">
var searchBox = new SearchBox("searchBox", "search",false,'Search');
</script>
<div id="navrow1" class="tabs">
<ul class="tablist">
<li><a href="index.html"><span>Main&#160;Page</span></a></li>
<li><a href="modules.html"><span>Modules</span></a></li>
<li><a href="namespaces.html"><span>Namespaces</span></a></li>
<li><a href="annotated.html"><span>Classes</span></a></li>
<li class="current"><a href="files.html"><span>Files</span></a></li>
<li>
<div id="MSearchBox" class="MSearchBoxInactive">
<span class="left">
<img id="MSearchSelect" src="search/mag_sel.png"
onmouseover="return searchBox.OnSearchSelectShow()"
onmouseout="return searchBox.OnSearchSelectHide()"
alt=""/>
<input type="text" id="MSearchField" value="Search" accesskey="S"
onfocus="searchBox.OnSearchFieldFocus(true)"
onblur="searchBox.OnSearchFieldFocus(false)"
onkeyup="searchBox.OnSearchFieldChange(event)"/>
</span><span class="right">
<a id="MSearchClose" href="javascript:searchBox.CloseResultsWindow()"><img id="MSearchCloseImg" border="0" src="search/close.png" alt=""/></a>
</span>
</div>
</li>
</ul>
</div>
<div id="navrow2" class="tabs2">
<ul class="tablist">
<li><a href="files.html"><span>File&#160;List</span></a></li>
<li><a href="globals.html"><span>File&#160;Members</span></a></li>
</ul>
</div>
<!-- window showing the filter options -->
<div id="MSearchSelectWindow"
onmouseover="return searchBox.OnSearchSelectShow()"
onmouseout="return searchBox.OnSearchSelectHide()"
onkeydown="return searchBox.OnSearchSelectKey(event)">
</div>
<!-- iframe showing the search results (closed by default) -->
<div id="MSearchResultsWindow">
<iframe src="javascript:void(0)" frameborder="0"
name="MSearchResults" id="MSearchResults">
</iframe>
</div>
<div id="nav-path" class="navpath">
<ul>
<li class="navelem"><a class="el" href="dir_d44c64559bbebec7f509842c48db8b23.html">include</a></li><li class="navelem"><a class="el" href="dir_6baf2bb612a2f0daa69af3101ede80a1.html">cutlass</a></li><li class="navelem"><a class="el" href="dir_048c1df36ab9c2efbb0733edba6291c9.html">arch</a></li> </ul>
</div>
</div><!-- top -->
<div class="header">
<div class="summary">
<a href="#nested-classes">Classes</a> &#124;
<a href="#namespaces">Namespaces</a> </div>
<div class="headertitle">
<div class="title">arch.h File Reference</div> </div>
</div><!--header-->
<div class="contents">
<p>Defines tags for architecture-specific configurations.
<a href="#details">More...</a></p>
<div class="textblock"><div class="dynheader">
This graph shows which files directly or indirectly include this file:</div>
<div class="dyncontent">
<div class="center"><img src="arch_8h__dep__incl.png" border="0" usemap="#arch_8hdep" alt=""/></div>
<map name="arch_8hdep" id="arch_8hdep">
</map>
</div>
</div>
<p><a href="arch_8h_source.html">Go to the source code of this file.</a></p>
<table class="memberdecls">
<tr class="heading"><td colspan="2"><h2 class="groupheader"><a name="nested-classes"></a>
Classes</h2></td></tr>
<tr class="memitem:"><td class="memItemLeft" align="right" valign="top">struct &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="structcutlass_1_1arch_1_1Sm50.html">cutlass::arch::Sm50</a></td></tr>
<tr class="separator:"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:"><td class="memItemLeft" align="right" valign="top">struct &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="structcutlass_1_1arch_1_1Sm60.html">cutlass::arch::Sm60</a></td></tr>
<tr class="separator:"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:"><td class="memItemLeft" align="right" valign="top">struct &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="structcutlass_1_1arch_1_1Sm61.html">cutlass::arch::Sm61</a></td></tr>
<tr class="separator:"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:"><td class="memItemLeft" align="right" valign="top">struct &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="structcutlass_1_1arch_1_1Sm70.html">cutlass::arch::Sm70</a></td></tr>
<tr class="separator:"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:"><td class="memItemLeft" align="right" valign="top">struct &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="structcutlass_1_1arch_1_1Sm72.html">cutlass::arch::Sm72</a></td></tr>
<tr class="separator:"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:"><td class="memItemLeft" align="right" valign="top">struct &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="structcutlass_1_1arch_1_1Sm75.html">cutlass::arch::Sm75</a></td></tr>
<tr class="separator:"><td class="memSeparator" colspan="2">&#160;</td></tr>
</table><table class="memberdecls">
<tr class="heading"><td colspan="2"><h2 class="groupheader"><a name="namespaces"></a>
Namespaces</h2></td></tr>
<tr class="memitem:namespacecutlass"><td class="memItemLeft" align="right" valign="top"> &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="namespacecutlass.html">cutlass</a></td></tr>
<tr class="separator:"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:namespacecutlass_1_1arch"><td class="memItemLeft" align="right" valign="top"> &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="namespacecutlass_1_1arch.html">cutlass::arch</a></td></tr>
<tr class="separator:"><td class="memSeparator" colspan="2">&#160;</td></tr>
</table>
</div><!-- contents -->
<!-- start footer part -->
<hr class="footer"/><address class="footer"><small>
Generated by &#160;<a href="http://www.doxygen.org/index.html">
<img class="footer" src="doxygen.png" alt="doxygen"/>
</a> 1.8.11
</small></address>
</body>
</html>

View File

@ -0,0 +1 @@
9ea32ea41ab87776449ab855965480b3

117
docs/arch_8h_source.html Normal file

File diff suppressed because one or more lines are too long

167
docs/array_8h.html Normal file
View File

@ -0,0 +1,167 @@
<!DOCTYPE html PUBLIC "-//W3C//DTD XHTML 1.0 Transitional//EN" "http://www.w3.org/TR/xhtml1/DTD/xhtml1-transitional.dtd">
<html xmlns="http://www.w3.org/1999/xhtml">
<head>
<meta http-equiv="Content-Type" content="text/xhtml;charset=UTF-8"/>
<meta http-equiv="X-UA-Compatible" content="IE=9"/>
<meta name="generator" content="Doxygen 1.8.11"/>
<title>CUTLASS: array.h File Reference</title>
<link href="tabs.css" rel="stylesheet" type="text/css"/>
<script type="text/javascript" src="jquery.js"></script>
<script type="text/javascript" src="dynsections.js"></script>
<link href="search/search.css" rel="stylesheet" type="text/css"/>
<script type="text/javascript" src="search/searchdata.js"></script>
<script type="text/javascript" src="search/search.js"></script>
<script type="text/javascript">
$(document).ready(function() { init_search(); });
</script>
<script type="text/x-mathjax-config">
MathJax.Hub.Config({
extensions: ["tex2jax.js"],
jax: ["input/TeX","output/HTML-CSS"],
});
</script><script type="text/javascript" src="http://cdn.mathjax.org/mathjax/latest/MathJax.js"></script>
<link href="doxygen.css" rel="stylesheet" type="text/css" />
</head>
<body>
<div id="top"><!-- do not remove this div, it is closed by doxygen! -->
<div id="titlearea">
<table cellspacing="0" cellpadding="0">
<tbody>
<tr style="height: 56px;">
<td id="projectlogo"><img alt="Logo" src="cutlass-logo-small.png"/></td>
<td id="projectalign" style="padding-left: 0.5em;">
<div id="projectname">CUTLASS
</div>
<div id="projectbrief">CUDA Templates for Linear Algebra Subroutines and Solvers</div>
</td>
</tr>
</tbody>
</table>
</div>
<!-- end header part -->
<!-- Generated by Doxygen 1.8.11 -->
<script type="text/javascript">
var searchBox = new SearchBox("searchBox", "search",false,'Search');
</script>
<div id="navrow1" class="tabs">
<ul class="tablist">
<li><a href="index.html"><span>Main&#160;Page</span></a></li>
<li><a href="modules.html"><span>Modules</span></a></li>
<li><a href="namespaces.html"><span>Namespaces</span></a></li>
<li><a href="annotated.html"><span>Classes</span></a></li>
<li class="current"><a href="files.html"><span>Files</span></a></li>
<li>
<div id="MSearchBox" class="MSearchBoxInactive">
<span class="left">
<img id="MSearchSelect" src="search/mag_sel.png"
onmouseover="return searchBox.OnSearchSelectShow()"
onmouseout="return searchBox.OnSearchSelectHide()"
alt=""/>
<input type="text" id="MSearchField" value="Search" accesskey="S"
onfocus="searchBox.OnSearchFieldFocus(true)"
onblur="searchBox.OnSearchFieldFocus(false)"
onkeyup="searchBox.OnSearchFieldChange(event)"/>
</span><span class="right">
<a id="MSearchClose" href="javascript:searchBox.CloseResultsWindow()"><img id="MSearchCloseImg" border="0" src="search/close.png" alt=""/></a>
</span>
</div>
</li>
</ul>
</div>
<div id="navrow2" class="tabs2">
<ul class="tablist">
<li><a href="files.html"><span>File&#160;List</span></a></li>
<li><a href="globals.html"><span>File&#160;Members</span></a></li>
</ul>
</div>
<!-- window showing the filter options -->
<div id="MSearchSelectWindow"
onmouseover="return searchBox.OnSearchSelectShow()"
onmouseout="return searchBox.OnSearchSelectHide()"
onkeydown="return searchBox.OnSearchSelectKey(event)">
</div>
<!-- iframe showing the search results (closed by default) -->
<div id="MSearchResultsWindow">
<iframe src="javascript:void(0)" frameborder="0"
name="MSearchResults" id="MSearchResults">
</iframe>
</div>
<div id="nav-path" class="navpath">
<ul>
<li class="navelem"><a class="el" href="dir_d44c64559bbebec7f509842c48db8b23.html">include</a></li><li class="navelem"><a class="el" href="dir_6baf2bb612a2f0daa69af3101ede80a1.html">cutlass</a></li> </ul>
</div>
</div><!-- top -->
<div class="header">
<div class="summary">
<a href="#nested-classes">Classes</a> &#124;
<a href="#namespaces">Namespaces</a> &#124;
<a href="#func-members">Functions</a> </div>
<div class="headertitle">
<div class="title">array.h File Reference</div> </div>
</div><!--header-->
<div class="contents">
<p>Statically sized array of elements that accommodates all CUTLASS-supported numeric types and is safe to use in a union.
<a href="#details">More...</a></p>
<div class="textblock"><code>#include &quot;<a class="el" href="cutlass_8h_source.html">cutlass/cutlass.h</a>&quot;</code><br />
<code>#include &quot;<a class="el" href="numeric__types_8h_source.html">cutlass/numeric_types.h</a>&quot;</code><br />
<code>#include &quot;<a class="el" href="array__subbyte_8h_source.html">cutlass/array_subbyte.h</a>&quot;</code><br />
</div><div class="textblock"><div class="dynheader">
Include dependency graph for array.h:</div>
<div class="dyncontent">
<div class="center"><img src="array_8h__incl.png" border="0" usemap="#array_8h" alt=""/></div>
<map name="array_8h" id="array_8h">
</map>
</div>
</div>
<p><a href="array_8h_source.html">Go to the source code of this file.</a></p>
<table class="memberdecls">
<tr class="heading"><td colspan="2"><h2 class="groupheader"><a name="nested-classes"></a>
Classes</h2></td></tr>
<tr class="memitem:"><td class="memItemLeft" align="right" valign="top">struct &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="structcutlass_1_1sizeof__bits_3_01Array_3_01T_00_01N_00_01RegisterSized_01_4_01_4.html">cutlass::sizeof_bits&lt; Array&lt; T, N, RegisterSized &gt; &gt;</a></td></tr>
<tr class="memdesc:"><td class="mdescLeft">&#160;</td><td class="mdescRight">Statically sized array for any data type. <a href="structcutlass_1_1sizeof__bits_3_01Array_3_01T_00_01N_00_01RegisterSized_01_4_01_4.html#details">More...</a><br /></td></tr>
<tr class="separator:"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:"><td class="memItemLeft" align="right" valign="top">class &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="classcutlass_1_1Array_3_01T_00_01N_00_01true_01_4.html">cutlass::Array&lt; T, N, true &gt;</a></td></tr>
<tr class="memdesc:"><td class="mdescLeft">&#160;</td><td class="mdescRight">Statically sized array for any data type. <a href="classcutlass_1_1Array_3_01T_00_01N_00_01true_01_4.html#details">More...</a><br /></td></tr>
<tr class="separator:"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:"><td class="memItemLeft" align="right" valign="top">class &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="classcutlass_1_1Array_3_01T_00_01N_00_01true_01_4_1_1iterator.html">cutlass::Array&lt; T, N, true &gt;::iterator</a></td></tr>
<tr class="memdesc:"><td class="mdescLeft">&#160;</td><td class="mdescRight">Bidirectional iterator over elements. <a href="classcutlass_1_1Array_3_01T_00_01N_00_01true_01_4_1_1iterator.html#details">More...</a><br /></td></tr>
<tr class="separator:"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:"><td class="memItemLeft" align="right" valign="top">class &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="classcutlass_1_1Array_3_01T_00_01N_00_01true_01_4_1_1const__iterator.html">cutlass::Array&lt; T, N, true &gt;::const_iterator</a></td></tr>
<tr class="memdesc:"><td class="mdescLeft">&#160;</td><td class="mdescRight">Bidirectional constant iterator over elements. <a href="classcutlass_1_1Array_3_01T_00_01N_00_01true_01_4_1_1const__iterator.html#details">More...</a><br /></td></tr>
<tr class="separator:"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:"><td class="memItemLeft" align="right" valign="top">class &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="classcutlass_1_1Array_3_01T_00_01N_00_01true_01_4_1_1reverse__iterator.html">cutlass::Array&lt; T, N, true &gt;::reverse_iterator</a></td></tr>
<tr class="memdesc:"><td class="mdescLeft">&#160;</td><td class="mdescRight">Bidirectional iterator over elements. <a href="classcutlass_1_1Array_3_01T_00_01N_00_01true_01_4_1_1reverse__iterator.html#details">More...</a><br /></td></tr>
<tr class="separator:"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:"><td class="memItemLeft" align="right" valign="top">class &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="classcutlass_1_1Array_3_01T_00_01N_00_01true_01_4_1_1const__reverse__iterator.html">cutlass::Array&lt; T, N, true &gt;::const_reverse_iterator</a></td></tr>
<tr class="memdesc:"><td class="mdescLeft">&#160;</td><td class="mdescRight">Bidirectional constant iterator over elements. <a href="classcutlass_1_1Array_3_01T_00_01N_00_01true_01_4_1_1const__reverse__iterator.html#details">More...</a><br /></td></tr>
<tr class="separator:"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:"><td class="memItemLeft" align="right" valign="top">class &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="classcutlass_1_1AlignedArray.html">cutlass::AlignedArray&lt; T, N, Alignment &gt;</a></td></tr>
<tr class="memdesc:"><td class="mdescLeft">&#160;</td><td class="mdescRight">Aligned array type. <a href="classcutlass_1_1AlignedArray.html#details">More...</a><br /></td></tr>
<tr class="separator:"><td class="memSeparator" colspan="2">&#160;</td></tr>
</table><table class="memberdecls">
<tr class="heading"><td colspan="2"><h2 class="groupheader"><a name="namespaces"></a>
Namespaces</h2></td></tr>
<tr class="memitem:namespacecutlass"><td class="memItemLeft" align="right" valign="top"> &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="namespacecutlass.html">cutlass</a></td></tr>
<tr class="separator:"><td class="memSeparator" colspan="2">&#160;</td></tr>
</table><table class="memberdecls">
<tr class="heading"><td colspan="2"><h2 class="groupheader"><a name="func-members"></a>
Functions</h2></td></tr>
<tr class="memitem:a935aabfdc47cf03f87c67bb22533f97f"><td class="memItemLeft" align="right" valign="top"><a class="el" href="cutlass_8h.html#a28c2443a142676d3d71effdae1a986b1">CUTLASS_HOST_DEVICE</a> <a class="el" href="platform_8h.html#a72f0657181cca64b44eb186b707eb380">constexpr</a> bool&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="namespacecutlass.html#a935aabfdc47cf03f87c67bb22533f97f">cutlass::ispow2</a> (unsigned x)</td></tr>
<tr class="memdesc:a935aabfdc47cf03f87c67bb22533f97f"><td class="mdescLeft">&#160;</td><td class="mdescRight">Returns true if the argument is a power of 2. <a href="namespacecutlass.html#a935aabfdc47cf03f87c67bb22533f97f">More...</a><br /></td></tr>
<tr class="separator:a935aabfdc47cf03f87c67bb22533f97f"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:ac16d8caf23537912eb02123c4bdacd14"><td class="memItemLeft" align="right" valign="top"><a class="el" href="cutlass_8h.html#a28c2443a142676d3d71effdae1a986b1">CUTLASS_HOST_DEVICE</a> <a class="el" href="platform_8h.html#a72f0657181cca64b44eb186b707eb380">constexpr</a> unsigned&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="namespacecutlass.html#ac16d8caf23537912eb02123c4bdacd14">cutlass::floor_pow_2</a> (unsigned x)</td></tr>
<tr class="memdesc:ac16d8caf23537912eb02123c4bdacd14"><td class="mdescLeft">&#160;</td><td class="mdescRight">Returns the largest power of two not greater than the argument. <a href="namespacecutlass.html#ac16d8caf23537912eb02123c4bdacd14">More...</a><br /></td></tr>
<tr class="separator:ac16d8caf23537912eb02123c4bdacd14"><td class="memSeparator" colspan="2">&#160;</td></tr>
</table>
</div><!-- contents -->
<!-- start footer part -->
<hr class="footer"/><address class="footer"><small>
Generated by &#160;<a href="http://www.doxygen.org/index.html">
<img class="footer" src="doxygen.png" alt="doxygen"/>
</a> 1.8.11
</small></address>
</body>
</html>

1
docs/array_8h__incl.md5 Normal file
View File

@ -0,0 +1 @@
90c159bd7ad938ad2d6e263ea8402fe7

194
docs/array_8h_source.html Normal file

File diff suppressed because one or more lines are too long

164
docs/array__subbyte_8h.html Normal file
View File

@ -0,0 +1,164 @@
<!DOCTYPE html PUBLIC "-//W3C//DTD XHTML 1.0 Transitional//EN" "http://www.w3.org/TR/xhtml1/DTD/xhtml1-transitional.dtd">
<html xmlns="http://www.w3.org/1999/xhtml">
<head>
<meta http-equiv="Content-Type" content="text/xhtml;charset=UTF-8"/>
<meta http-equiv="X-UA-Compatible" content="IE=9"/>
<meta name="generator" content="Doxygen 1.8.11"/>
<title>CUTLASS: array_subbyte.h File Reference</title>
<link href="tabs.css" rel="stylesheet" type="text/css"/>
<script type="text/javascript" src="jquery.js"></script>
<script type="text/javascript" src="dynsections.js"></script>
<link href="search/search.css" rel="stylesheet" type="text/css"/>
<script type="text/javascript" src="search/searchdata.js"></script>
<script type="text/javascript" src="search/search.js"></script>
<script type="text/javascript">
$(document).ready(function() { init_search(); });
</script>
<script type="text/x-mathjax-config">
MathJax.Hub.Config({
extensions: ["tex2jax.js"],
jax: ["input/TeX","output/HTML-CSS"],
});
</script><script type="text/javascript" src="http://cdn.mathjax.org/mathjax/latest/MathJax.js"></script>
<link href="doxygen.css" rel="stylesheet" type="text/css" />
</head>
<body>
<div id="top"><!-- do not remove this div, it is closed by doxygen! -->
<div id="titlearea">
<table cellspacing="0" cellpadding="0">
<tbody>
<tr style="height: 56px;">
<td id="projectlogo"><img alt="Logo" src="cutlass-logo-small.png"/></td>
<td id="projectalign" style="padding-left: 0.5em;">
<div id="projectname">CUTLASS
</div>
<div id="projectbrief">CUDA Templates for Linear Algebra Subroutines and Solvers</div>
</td>
</tr>
</tbody>
</table>
</div>
<!-- end header part -->
<!-- Generated by Doxygen 1.8.11 -->
<script type="text/javascript">
var searchBox = new SearchBox("searchBox", "search",false,'Search');
</script>
<div id="navrow1" class="tabs">
<ul class="tablist">
<li><a href="index.html"><span>Main&#160;Page</span></a></li>
<li><a href="modules.html"><span>Modules</span></a></li>
<li><a href="namespaces.html"><span>Namespaces</span></a></li>
<li><a href="annotated.html"><span>Classes</span></a></li>
<li class="current"><a href="files.html"><span>Files</span></a></li>
<li>
<div id="MSearchBox" class="MSearchBoxInactive">
<span class="left">
<img id="MSearchSelect" src="search/mag_sel.png"
onmouseover="return searchBox.OnSearchSelectShow()"
onmouseout="return searchBox.OnSearchSelectHide()"
alt=""/>
<input type="text" id="MSearchField" value="Search" accesskey="S"
onfocus="searchBox.OnSearchFieldFocus(true)"
onblur="searchBox.OnSearchFieldFocus(false)"
onkeyup="searchBox.OnSearchFieldChange(event)"/>
</span><span class="right">
<a id="MSearchClose" href="javascript:searchBox.CloseResultsWindow()"><img id="MSearchCloseImg" border="0" src="search/close.png" alt=""/></a>
</span>
</div>
</li>
</ul>
</div>
<div id="navrow2" class="tabs2">
<ul class="tablist">
<li><a href="files.html"><span>File&#160;List</span></a></li>
<li><a href="globals.html"><span>File&#160;Members</span></a></li>
</ul>
</div>
<!-- window showing the filter options -->
<div id="MSearchSelectWindow"
onmouseover="return searchBox.OnSearchSelectShow()"
onmouseout="return searchBox.OnSearchSelectHide()"
onkeydown="return searchBox.OnSearchSelectKey(event)">
</div>
<!-- iframe showing the search results (closed by default) -->
<div id="MSearchResultsWindow">
<iframe src="javascript:void(0)" frameborder="0"
name="MSearchResults" id="MSearchResults">
</iframe>
</div>
<div id="nav-path" class="navpath">
<ul>
<li class="navelem"><a class="el" href="dir_d44c64559bbebec7f509842c48db8b23.html">include</a></li><li class="navelem"><a class="el" href="dir_6baf2bb612a2f0daa69af3101ede80a1.html">cutlass</a></li> </ul>
</div>
</div><!-- top -->
<div class="header">
<div class="summary">
<a href="#nested-classes">Classes</a> &#124;
<a href="#namespaces">Namespaces</a> </div>
<div class="headertitle">
<div class="title">array_subbyte.h File Reference</div> </div>
</div><!--header-->
<div class="contents">
<p>Statically sized array of elements that accommodates all CUTLASS-supported numeric types and is safe to use in a union.
<a href="#details">More...</a></p>
<div class="textblock"><code>#include &quot;<a class="el" href="cutlass_8h_source.html">cutlass/cutlass.h</a>&quot;</code><br />
<code>#include &quot;<a class="el" href="array_8h_source.html">cutlass/array.h</a>&quot;</code><br />
<code>#include &quot;<a class="el" href="platform_8h_source.html">cutlass/platform/platform.h</a>&quot;</code><br />
</div><div class="textblock"><div class="dynheader">
Include dependency graph for array_subbyte.h:</div>
<div class="dyncontent">
<div class="center"><img src="array__subbyte_8h__incl.png" border="0" usemap="#array__subbyte_8h" alt=""/></div>
<map name="array__subbyte_8h" id="array__subbyte_8h">
</map>
</div>
</div><div class="textblock"><div class="dynheader">
This graph shows which files directly or indirectly include this file:</div>
<div class="dyncontent">
<div class="center"><img src="array__subbyte_8h__dep__incl.png" border="0" usemap="#array__subbyte_8hdep" alt=""/></div>
<map name="array__subbyte_8hdep" id="array__subbyte_8hdep">
</map>
</div>
</div>
<p><a href="array__subbyte_8h_source.html">Go to the source code of this file.</a></p>
<table class="memberdecls">
<tr class="heading"><td colspan="2"><h2 class="groupheader"><a name="nested-classes"></a>
Classes</h2></td></tr>
<tr class="memitem:"><td class="memItemLeft" align="right" valign="top">class &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html">cutlass::Array&lt; T, N, false &gt;</a></td></tr>
<tr class="memdesc:"><td class="mdescLeft">&#160;</td><td class="mdescRight">Statically sized array for any data type. <a href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4.html#details">More...</a><br /></td></tr>
<tr class="separator:"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:"><td class="memItemLeft" align="right" valign="top">class &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1reference.html">cutlass::Array&lt; T, N, false &gt;::reference</a></td></tr>
<tr class="memdesc:"><td class="mdescLeft">&#160;</td><td class="mdescRight">Reference object inserts or extracts sub-byte items. <a href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1reference.html#details">More...</a><br /></td></tr>
<tr class="separator:"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:"><td class="memItemLeft" align="right" valign="top">class &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1const__reference.html">cutlass::Array&lt; T, N, false &gt;::const_reference</a></td></tr>
<tr class="memdesc:"><td class="mdescLeft">&#160;</td><td class="mdescRight">Reference object extracts sub-byte items. <a href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1const__reference.html#details">More...</a><br /></td></tr>
<tr class="separator:"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:"><td class="memItemLeft" align="right" valign="top">class &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1iterator.html">cutlass::Array&lt; T, N, false &gt;::iterator</a></td></tr>
<tr class="memdesc:"><td class="mdescLeft">&#160;</td><td class="mdescRight">Bidirectional iterator over elements. <a href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1iterator.html#details">More...</a><br /></td></tr>
<tr class="separator:"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:"><td class="memItemLeft" align="right" valign="top">class &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1const__iterator.html">cutlass::Array&lt; T, N, false &gt;::const_iterator</a></td></tr>
<tr class="memdesc:"><td class="mdescLeft">&#160;</td><td class="mdescRight">Bidirectional constant iterator over elements. <a href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1const__iterator.html#details">More...</a><br /></td></tr>
<tr class="separator:"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:"><td class="memItemLeft" align="right" valign="top">class &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1reverse__iterator.html">cutlass::Array&lt; T, N, false &gt;::reverse_iterator</a></td></tr>
<tr class="memdesc:"><td class="mdescLeft">&#160;</td><td class="mdescRight">Bidirectional iterator over elements. <a href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1reverse__iterator.html#details">More...</a><br /></td></tr>
<tr class="separator:"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:"><td class="memItemLeft" align="right" valign="top">class &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1const__reverse__iterator.html">cutlass::Array&lt; T, N, false &gt;::const_reverse_iterator</a></td></tr>
<tr class="memdesc:"><td class="mdescLeft">&#160;</td><td class="mdescRight">Bidirectional constant iterator over elements. <a href="classcutlass_1_1Array_3_01T_00_01N_00_01false_01_4_1_1const__reverse__iterator.html#details">More...</a><br /></td></tr>
<tr class="separator:"><td class="memSeparator" colspan="2">&#160;</td></tr>
</table><table class="memberdecls">
<tr class="heading"><td colspan="2"><h2 class="groupheader"><a name="namespaces"></a>
Namespaces</h2></td></tr>
<tr class="memitem:namespacecutlass"><td class="memItemLeft" align="right" valign="top"> &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="namespacecutlass.html">cutlass</a></td></tr>
<tr class="separator:"><td class="memSeparator" colspan="2">&#160;</td></tr>
</table>
</div><!-- contents -->
<!-- start footer part -->
<hr class="footer"/><address class="footer"><small>
Generated by &#160;<a href="http://www.doxygen.org/index.html">
<img class="footer" src="doxygen.png" alt="doxygen"/>
</a> 1.8.11
</small></address>
</body>
</html>

View File

@ -0,0 +1 @@
7c0288c037b6ea169ec7a3aa1015a4d4

View File

@ -0,0 +1 @@
36310516438810c2a8ba31a7816cd1de

File diff suppressed because one or more lines are too long

View File

@ -0,0 +1,155 @@
<!DOCTYPE html PUBLIC "-//W3C//DTD XHTML 1.0 Transitional//EN" "http://www.w3.org/TR/xhtml1/DTD/xhtml1-transitional.dtd">
<html xmlns="http://www.w3.org/1999/xhtml">
<head>
<meta http-equiv="Content-Type" content="text/xhtml;charset=UTF-8"/>
<meta http-equiv="X-UA-Compatible" content="IE=9"/>
<meta name="generator" content="Doxygen 1.8.11"/>
<title>CUTLASS: batched_reduction.h File Reference</title>
<link href="tabs.css" rel="stylesheet" type="text/css"/>
<script type="text/javascript" src="jquery.js"></script>
<script type="text/javascript" src="dynsections.js"></script>
<link href="search/search.css" rel="stylesheet" type="text/css"/>
<script type="text/javascript" src="search/searchdata.js"></script>
<script type="text/javascript" src="search/search.js"></script>
<script type="text/javascript">
$(document).ready(function() { init_search(); });
</script>
<script type="text/x-mathjax-config">
MathJax.Hub.Config({
extensions: ["tex2jax.js"],
jax: ["input/TeX","output/HTML-CSS"],
});
</script><script type="text/javascript" src="http://cdn.mathjax.org/mathjax/latest/MathJax.js"></script>
<link href="doxygen.css" rel="stylesheet" type="text/css" />
</head>
<body>
<div id="top"><!-- do not remove this div, it is closed by doxygen! -->
<div id="titlearea">
<table cellspacing="0" cellpadding="0">
<tbody>
<tr style="height: 56px;">
<td id="projectlogo"><img alt="Logo" src="cutlass-logo-small.png"/></td>
<td id="projectalign" style="padding-left: 0.5em;">
<div id="projectname">CUTLASS
</div>
<div id="projectbrief">CUDA Templates for Linear Algebra Subroutines and Solvers</div>
</td>
</tr>
</tbody>
</table>
</div>
<!-- end header part -->
<!-- Generated by Doxygen 1.8.11 -->
<script type="text/javascript">
var searchBox = new SearchBox("searchBox", "search",false,'Search');
</script>
<div id="navrow1" class="tabs">
<ul class="tablist">
<li><a href="index.html"><span>Main&#160;Page</span></a></li>
<li><a href="modules.html"><span>Modules</span></a></li>
<li><a href="namespaces.html"><span>Namespaces</span></a></li>
<li><a href="annotated.html"><span>Classes</span></a></li>
<li class="current"><a href="files.html"><span>Files</span></a></li>
<li>
<div id="MSearchBox" class="MSearchBoxInactive">
<span class="left">
<img id="MSearchSelect" src="search/mag_sel.png"
onmouseover="return searchBox.OnSearchSelectShow()"
onmouseout="return searchBox.OnSearchSelectHide()"
alt=""/>
<input type="text" id="MSearchField" value="Search" accesskey="S"
onfocus="searchBox.OnSearchFieldFocus(true)"
onblur="searchBox.OnSearchFieldFocus(false)"
onkeyup="searchBox.OnSearchFieldChange(event)"/>
</span><span class="right">
<a id="MSearchClose" href="javascript:searchBox.CloseResultsWindow()"><img id="MSearchCloseImg" border="0" src="search/close.png" alt=""/></a>
</span>
</div>
</li>
</ul>
</div>
<div id="navrow2" class="tabs2">
<ul class="tablist">
<li><a href="files.html"><span>File&#160;List</span></a></li>
<li><a href="globals.html"><span>File&#160;Members</span></a></li>
</ul>
</div>
<!-- window showing the filter options -->
<div id="MSearchSelectWindow"
onmouseover="return searchBox.OnSearchSelectShow()"
onmouseout="return searchBox.OnSearchSelectHide()"
onkeydown="return searchBox.OnSearchSelectKey(event)">
</div>
<!-- iframe showing the search results (closed by default) -->
<div id="MSearchResultsWindow">
<iframe src="javascript:void(0)" frameborder="0"
name="MSearchResults" id="MSearchResults">
</iframe>
</div>
<div id="nav-path" class="navpath">
<ul>
<li class="navelem"><a class="el" href="dir_d44c64559bbebec7f509842c48db8b23.html">include</a></li><li class="navelem"><a class="el" href="dir_6baf2bb612a2f0daa69af3101ede80a1.html">cutlass</a></li><li class="navelem"><a class="el" href="dir_ac488927e63b76ba9cb3ad9c317bbde9.html">reduction</a></li> </ul>
</div>
</div><!-- top -->
<div class="header">
<div class="summary">
<a href="#nested-classes">Classes</a> &#124;
<a href="#namespaces">Namespaces</a> &#124;
<a href="#func-members">Functions</a> </div>
<div class="headertitle">
<div class="title">batched_reduction.h File Reference</div> </div>
</div><!--header-->
<div class="contents">
<p>Implements a software-pipelined efficient batched reduction. D = alpha * Reduction(A) + beta * C.
<a href="#details">More...</a></p>
<div class="textblock"><code>#include &lt;cuda.h&gt;</code><br />
<code>#include &quot;<a class="el" href="coord_8h_source.html">cutlass/coord.h</a>&quot;</code><br />
<code>#include &quot;cutlass/util/platform.h&quot;</code><br />
<code>#include &quot;cutlass/fragment.h&quot;</code><br />
</div><div class="textblock"><div class="dynheader">
Include dependency graph for batched_reduction.h:</div>
<div class="dyncontent">
<div class="center"><img src="batched__reduction_8h__incl.png" border="0" usemap="#batched__reduction_8h" alt=""/></div>
<map name="batched__reduction_8h" id="batched__reduction_8h">
</map>
</div>
</div><div class="textblock"><div class="dynheader">
This graph shows which files directly or indirectly include this file:</div>
<div class="dyncontent">
<div class="center"><img src="batched__reduction_8h__dep__incl.png" border="0" usemap="#batched__reduction_8hdep" alt=""/></div>
<map name="batched__reduction_8hdep" id="batched__reduction_8hdep">
</map>
</div>
</div>
<p><a href="batched__reduction_8h_source.html">Go to the source code of this file.</a></p>
<table class="memberdecls">
<tr class="heading"><td colspan="2"><h2 class="groupheader"><a name="nested-classes"></a>
Classes</h2></td></tr>
<tr class="memitem:"><td class="memItemLeft" align="right" valign="top">struct &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="structcutlass_1_1reduction_1_1BatchedReduction.html">cutlass::reduction::BatchedReduction&lt; BatchedReductionTraits_ &gt;</a></td></tr>
<tr class="separator:"><td class="memSeparator" colspan="2">&#160;</td></tr>
</table><table class="memberdecls">
<tr class="heading"><td colspan="2"><h2 class="groupheader"><a name="namespaces"></a>
Namespaces</h2></td></tr>
<tr class="memitem:namespacecutlass"><td class="memItemLeft" align="right" valign="top"> &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="namespacecutlass.html">cutlass</a></td></tr>
<tr class="separator:"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:namespacecutlass_1_1reduction"><td class="memItemLeft" align="right" valign="top"> &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="namespacecutlass_1_1reduction.html">cutlass::reduction</a></td></tr>
<tr class="separator:"><td class="memSeparator" colspan="2">&#160;</td></tr>
</table><table class="memberdecls">
<tr class="heading"><td colspan="2"><h2 class="groupheader"><a name="func-members"></a>
Functions</h2></td></tr>
<tr class="memitem:a9665e8f438a7b290d6e2eb640d93045f"><td class="memTemplParams" colspan="2">template&lt;typename batched_reduction_ &gt; </td></tr>
<tr class="memitem:a9665e8f438a7b290d6e2eb640d93045f"><td class="memTemplItemLeft" align="right" valign="top">__global__&#160;</td><td class="memTemplItemRight" valign="bottom"><a class="el" href="namespacecutlass_1_1reduction.html#a9665e8f438a7b290d6e2eb640d93045f">cutlass::reduction::__launch_bounds__</a> (batched_reduction_::Traits::kThreads, 1) void batched_reduction_kernel(typename batched_reduction_</td></tr>
<tr class="separator:a9665e8f438a7b290d6e2eb640d93045f"><td class="memSeparator" colspan="2">&#160;</td></tr>
</table>
</div><!-- contents -->
<!-- start footer part -->
<hr class="footer"/><address class="footer"><small>
Generated by &#160;<a href="http://www.doxygen.org/index.html">
<img class="footer" src="doxygen.png" alt="doxygen"/>
</a> 1.8.11
</small></address>
</body>
</html>

View File

@ -0,0 +1 @@
2bce650f452329d669d303788cc619c8

View File

@ -0,0 +1 @@
d38876c9b9d3ade81fb457e3ebf5c6fd

File diff suppressed because one or more lines are too long

View File

@ -0,0 +1,144 @@
<!DOCTYPE html PUBLIC "-//W3C//DTD XHTML 1.0 Transitional//EN" "http://www.w3.org/TR/xhtml1/DTD/xhtml1-transitional.dtd">
<html xmlns="http://www.w3.org/1999/xhtml">
<head>
<meta http-equiv="Content-Type" content="text/xhtml;charset=UTF-8"/>
<meta http-equiv="X-UA-Compatible" content="IE=9"/>
<meta name="generator" content="Doxygen 1.8.11"/>
<title>CUTLASS: batched_reduction_traits.h File Reference</title>
<link href="tabs.css" rel="stylesheet" type="text/css"/>
<script type="text/javascript" src="jquery.js"></script>
<script type="text/javascript" src="dynsections.js"></script>
<link href="search/search.css" rel="stylesheet" type="text/css"/>
<script type="text/javascript" src="search/searchdata.js"></script>
<script type="text/javascript" src="search/search.js"></script>
<script type="text/javascript">
$(document).ready(function() { init_search(); });
</script>
<script type="text/x-mathjax-config">
MathJax.Hub.Config({
extensions: ["tex2jax.js"],
jax: ["input/TeX","output/HTML-CSS"],
});
</script><script type="text/javascript" src="http://cdn.mathjax.org/mathjax/latest/MathJax.js"></script>
<link href="doxygen.css" rel="stylesheet" type="text/css" />
</head>
<body>
<div id="top"><!-- do not remove this div, it is closed by doxygen! -->
<div id="titlearea">
<table cellspacing="0" cellpadding="0">
<tbody>
<tr style="height: 56px;">
<td id="projectlogo"><img alt="Logo" src="cutlass-logo-small.png"/></td>
<td id="projectalign" style="padding-left: 0.5em;">
<div id="projectname">CUTLASS
</div>
<div id="projectbrief">CUDA Templates for Linear Algebra Subroutines and Solvers</div>
</td>
</tr>
</tbody>
</table>
</div>
<!-- end header part -->
<!-- Generated by Doxygen 1.8.11 -->
<script type="text/javascript">
var searchBox = new SearchBox("searchBox", "search",false,'Search');
</script>
<div id="navrow1" class="tabs">
<ul class="tablist">
<li><a href="index.html"><span>Main&#160;Page</span></a></li>
<li><a href="modules.html"><span>Modules</span></a></li>
<li><a href="namespaces.html"><span>Namespaces</span></a></li>
<li><a href="annotated.html"><span>Classes</span></a></li>
<li class="current"><a href="files.html"><span>Files</span></a></li>
<li>
<div id="MSearchBox" class="MSearchBoxInactive">
<span class="left">
<img id="MSearchSelect" src="search/mag_sel.png"
onmouseover="return searchBox.OnSearchSelectShow()"
onmouseout="return searchBox.OnSearchSelectHide()"
alt=""/>
<input type="text" id="MSearchField" value="Search" accesskey="S"
onfocus="searchBox.OnSearchFieldFocus(true)"
onblur="searchBox.OnSearchFieldFocus(false)"
onkeyup="searchBox.OnSearchFieldChange(event)"/>
</span><span class="right">
<a id="MSearchClose" href="javascript:searchBox.CloseResultsWindow()"><img id="MSearchCloseImg" border="0" src="search/close.png" alt=""/></a>
</span>
</div>
</li>
</ul>
</div>
<div id="navrow2" class="tabs2">
<ul class="tablist">
<li><a href="files.html"><span>File&#160;List</span></a></li>
<li><a href="globals.html"><span>File&#160;Members</span></a></li>
</ul>
</div>
<!-- window showing the filter options -->
<div id="MSearchSelectWindow"
onmouseover="return searchBox.OnSearchSelectShow()"
onmouseout="return searchBox.OnSearchSelectHide()"
onkeydown="return searchBox.OnSearchSelectKey(event)">
</div>
<!-- iframe showing the search results (closed by default) -->
<div id="MSearchResultsWindow">
<iframe src="javascript:void(0)" frameborder="0"
name="MSearchResults" id="MSearchResults">
</iframe>
</div>
<div id="nav-path" class="navpath">
<ul>
<li class="navelem"><a class="el" href="dir_d44c64559bbebec7f509842c48db8b23.html">include</a></li><li class="navelem"><a class="el" href="dir_6baf2bb612a2f0daa69af3101ede80a1.html">cutlass</a></li><li class="navelem"><a class="el" href="dir_ac488927e63b76ba9cb3ad9c317bbde9.html">reduction</a></li> </ul>
</div>
</div><!-- top -->
<div class="header">
<div class="summary">
<a href="#nested-classes">Classes</a> &#124;
<a href="#namespaces">Namespaces</a> </div>
<div class="headertitle">
<div class="title">batched_reduction_traits.h File Reference</div> </div>
</div><!--header-->
<div class="contents">
<p>Defines structural properties of complete batched reduction. D = alpha * Reduction(A) + beta * C.
<a href="#details">More...</a></p>
<div class="textblock"><code>#include &quot;<a class="el" href="cutlass_8h_source.html">cutlass/cutlass.h</a>&quot;</code><br />
<code>#include &quot;cutlass/shape.h&quot;</code><br />
<code>#include &quot;<a class="el" href="reduction_2threadblock__swizzle_8h_source.html">cutlass/reduction/threadblock_swizzle.h</a>&quot;</code><br />
<code>#include &quot;<a class="el" href="batched__reduction_8h_source.html">cutlass/reduction/batched_reduction.h</a>&quot;</code><br />
<code>#include &quot;cutlass/gemm/linear_scaling.h&quot;</code><br />
</div><div class="textblock"><div class="dynheader">
Include dependency graph for batched_reduction_traits.h:</div>
<div class="dyncontent">
<div class="center"><img src="batched__reduction__traits_8h__incl.png" border="0" usemap="#batched__reduction__traits_8h" alt=""/></div>
<map name="batched__reduction__traits_8h" id="batched__reduction__traits_8h">
</map>
</div>
</div>
<p><a href="batched__reduction__traits_8h_source.html">Go to the source code of this file.</a></p>
<table class="memberdecls">
<tr class="heading"><td colspan="2"><h2 class="groupheader"><a name="nested-classes"></a>
Classes</h2></td></tr>
<tr class="memitem:"><td class="memItemLeft" align="right" valign="top">struct &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="structcutlass_1_1reduction_1_1BatchedReductionTraits.html">cutlass::reduction::BatchedReductionTraits&lt; ScalarA_, ScalarC_, ScalarD_, ScalarAlphaBeta_, ScalarAccum_, ReductionSize_, OutputTile_, SubTile_, ThreadShape_, Index_, BlockSwizzle_, maxInReg_, maxOutReg_, Functor_ &gt;</a></td></tr>
<tr class="separator:"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:"><td class="memItemLeft" align="right" valign="top">struct &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="structcutlass_1_1reduction_1_1BatchedReductionTraits_1_1Params.html">cutlass::reduction::BatchedReductionTraits&lt; ScalarA_, ScalarC_, ScalarD_, ScalarAlphaBeta_, ScalarAccum_, ReductionSize_, OutputTile_, SubTile_, ThreadShape_, Index_, BlockSwizzle_, maxInReg_, maxOutReg_, Functor_ &gt;::Params</a></td></tr>
<tr class="separator:"><td class="memSeparator" colspan="2">&#160;</td></tr>
</table><table class="memberdecls">
<tr class="heading"><td colspan="2"><h2 class="groupheader"><a name="namespaces"></a>
Namespaces</h2></td></tr>
<tr class="memitem:namespacecutlass"><td class="memItemLeft" align="right" valign="top"> &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="namespacecutlass.html">cutlass</a></td></tr>
<tr class="separator:"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:namespacecutlass_1_1reduction"><td class="memItemLeft" align="right" valign="top"> &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="namespacecutlass_1_1reduction.html">cutlass::reduction</a></td></tr>
<tr class="separator:"><td class="memSeparator" colspan="2">&#160;</td></tr>
</table>
</div><!-- contents -->
<!-- start footer part -->
<hr class="footer"/><address class="footer"><small>
Generated by &#160;<a href="http://www.doxygen.org/index.html">
<img class="footer" src="doxygen.png" alt="doxygen"/>
</a> 1.8.11
</small></address>
</body>
</html>

View File

@ -0,0 +1 @@
957af6c3e40d98d122a3ef83474f7252

Some files were not shown because too many files have changed in this diff Show More