From d7137f9c0a1633b76455109373887e1640713b5d Mon Sep 17 00:00:00 2001 From: akerr Date: Wed, 19 Sep 2018 14:02:08 -0700 Subject: [PATCH] Updated doxygen --- docs/annotated.html | 649 ++++---- ...dicateVector_1_1ConstIterator-members.html | 2 +- ...s_1_1PredicateVector_1_1ConstIterator.html | 2 +- ..._1PredicateVector_1_1Iterator-members.html | 2 +- ...utlass_1_1PredicateVector_1_1Iterator.html | 2 +- docs/classcutlass_1_1TensorRef-members.html | 61 +- docs/classcutlass_1_1TensorRef.html | 1318 ++++++++++------ docs/classcutlass_1_1TensorRef.png | Bin 0 -> 2451 bytes ...nsorRefArray_1_1ConstIterator-members.html | 101 ++ ...ss_1_1TensorRefArray_1_1ConstIterator.html | 440 ++++++ ...BatchStrided_1_1ConstIterator-members.html | 102 ++ ...ensorRefBatchStrided_1_1ConstIterator.html | 476 ++++++ ...Index___00_01LongIndex___01_4-members.html | 124 ++ ...11_00_01Index___00_01LongIndex___01_4.html | 1092 +++++++++++++ docs/classcutlass_1_1TensorView-members.html | 95 +- docs/classcutlass_1_1TensorView.html | 1289 +++++++++------- docs/classcutlass_1_1TensorView.png | Bin 690 -> 1605 bytes ...asscutlass_1_1ZipTileIterator-members.html | 125 ++ docs/classcutlass_1_1ZipTileIterator.html | 1290 ++++++++++++++++ ..._1_1detail_1_1ScalarOrPointer-members.html | 101 ++ ...scutlass_1_1detail_1_1ScalarOrPointer.html | 434 ++++++ ...earScalingDevicePtr_1_1Params-members.html | 98 ++ ...m_1_1LinearScalingDevicePtr_1_1Params.html | 389 +++++ ...utlass_1_1platform_1_1complex-members.html | 100 ++ docs/classcutlass_1_1platform_1_1complex.html | 413 +++++ ...ss_1_1platform_1_1unique__ptr-members.html | 2 +- ...asscutlass_1_1platform_1_1unique__ptr.html | 2 +- docs/classes.html | 174 ++- docs/clear__accumulators_8h.html | 4 +- docs/clear__accumulators_8h_source.html | 7 +- docs/complex_8h.html | 263 ++++ docs/complex_8h_source.html | 123 ++ docs/convert_8h.html | 4 +- docs/convert_8h_source.html | 4 +- docs/coord_8h.html | 22 +- docs/coord_8h_source.html | 80 +- docs/core__io_8h.html | 79 +- docs/core__io_8h_source.html | 12 +- docs/cutlass_8h.html | 64 +- docs/cutlass_8h_source.html | 7 +- docs/cutlass__math_8h.html | 14 +- docs/cutlass__math_8h_source.html | 14 +- docs/debug_8h.html | 2 +- docs/debug_8h_source.html | 2 +- docs/dgemm__traits_8h.html | 20 +- docs/dgemm__traits_8h_source.html | 16 +- .../dir_1417ee5ebebc309c36b7962f26a92c39.html | 29 +- .../dir_18d6a367a3982a494d65599933fc67a3.html | 32 +- .../dir_c5917a9a879e9a6c73eaf5237444ab84.html | 6 +- docs/files.html | 119 +- docs/fp16__sgemm__multiply__add_8h.html | 111 ++ .../fp16__sgemm__multiply__add_8h_source.html | 107 ++ docs/fp16__sgemm__traits_8h.html | 117 ++ docs/fp16__sgemm__traits_8h_source.html | 104 ++ docs/fragment_8h.html | 10 +- docs/fragment_8h_source.html | 82 +- docs/fragment__multiply__add_8h.html | 8 +- docs/fragment__multiply__add_8h_source.html | 30 +- docs/functions.html | 104 +- docs/functions_0x7e.html | 2 +- docs/functions_b.html | 68 +- docs/functions_c.html | 91 +- docs/functions_d.html | 60 +- docs/functions_e.html | 12 +- docs/functions_enum.html | 8 +- docs/functions_eval.html | 40 +- docs/functions_f.html | 132 +- docs/functions_func.html | 27 +- docs/functions_func_0x7e.html | 2 +- docs/functions_func_b.html | 10 +- docs/functions_func_c.html | 65 +- docs/functions_func_d.html | 24 +- docs/functions_func_e.html | 8 +- docs/functions_func_f.html | 19 +- docs/functions_func_g.html | 60 +- docs/functions_func_h.html | 11 +- docs/functions_func_i.html | 103 +- docs/functions_func_k.html | 98 ++ docs/functions_func_l.html | 50 +- docs/functions_func_m.html | 31 +- docs/functions_func_n.html | 89 ++ docs/functions_func_o.html | 151 +- docs/functions_func_p.html | 22 +- docs/functions_func_r.html | 35 +- docs/functions_func_s.html | 68 +- docs/functions_func_t.html | 40 +- docs/functions_func_u.html | 2 +- docs/functions_func_v.html | 18 +- docs/functions_func_w.html | 7 +- docs/functions_func_z.html | 95 ++ docs/functions_g.html | 133 +- docs/functions_h.html | 11 +- docs/functions_i.html | 214 +-- docs/functions_k.html | 203 ++- docs/functions_l.html | 77 +- docs/functions_m.html | 53 +- docs/functions_n.html | 13 +- docs/functions_o.html | 160 +- docs/functions_p.html | 89 +- docs/functions_r.html | 43 +- docs/functions_s.html | 374 +++-- docs/functions_t.html | 194 ++- docs/functions_type.html | 69 +- docs/functions_type_b.html | 35 +- docs/functions_type_c.html | 25 +- docs/functions_type_d.html | 24 +- docs/functions_type_e.html | 6 +- docs/functions_type_f.html | 64 +- docs/functions_type_g.html | 67 +- docs/functions_type_i.html | 75 +- docs/functions_type_k.html | 86 ++ docs/functions_type_l.html | 10 +- docs/functions_type_m.html | 15 +- docs/functions_type_n.html | 2 +- docs/functions_type_o.html | 9 +- docs/functions_type_p.html | 23 +- docs/functions_type_s.html | 211 +-- docs/functions_type_t.html | 117 +- docs/functions_type_v.html | 8 +- docs/functions_type_w.html | 4 +- docs/functions_type_y.html | 2 +- docs/functions_u.html | 2 +- docs/functions_v.html | 27 +- docs/functions_vars.html | 12 +- docs/functions_vars_b.html | 26 +- docs/functions_vars_c.html | 5 +- docs/functions_vars_d.html | 16 +- docs/functions_vars_e.html | 2 +- docs/functions_vars_f.html | 23 +- docs/functions_vars_g.html | 14 +- docs/functions_vars_i.html | 22 +- docs/functions_vars_k.html | 157 +- docs/functions_vars_l.html | 21 +- docs/functions_vars_m.html | 13 +- docs/functions_vars_n.html | 9 +- docs/functions_vars_o.html | 86 ++ docs/functions_vars_p.html | 28 +- docs/functions_vars_r.html | 10 +- docs/functions_vars_s.html | 73 +- docs/functions_vars_t.html | 29 +- docs/functions_vars_v.html | 3 +- docs/functions_w.html | 9 +- docs/functions_y.html | 2 +- docs/functions_z.html | 95 ++ docs/gemm_8h.html | 26 +- docs/gemm_8h_source.html | 74 +- docs/gemm__config_8h.html | 109 ++ docs/gemm__config_8h_source.html | 123 ++ docs/gemm__coord_8h.html | 111 ++ docs/gemm__coord_8h_source.html | 131 ++ docs/gemm__desc_8h.html | 111 ++ docs/gemm__desc_8h_source.html | 121 ++ docs/gemm__epilogue_8h.html | 19 +- docs/gemm__epilogue_8h_source.html | 73 +- docs/gemm__epilogue__traits_8h.html | 24 +- docs/gemm__epilogue__traits_8h_source.html | 121 +- docs/gemm__global__stream_8h.html | 20 +- docs/gemm__global__stream_8h_source.html | 93 +- docs/gemm__global__tile_8h.html | 16 +- docs/gemm__global__tile_8h_source.html | 235 +-- docs/gemm__operand_8h.html | 8 +- docs/gemm__operand_8h_source.html | 24 +- docs/gemm__shared__stream_8h.html | 5 +- docs/gemm__shared__stream_8h_source.html | 53 +- docs/gemm__shared__tile_8h.html | 4 +- docs/gemm__shared__tile_8h_source.html | 181 +-- docs/gemm__stream__pair_8h.html | 132 ++ docs/gemm__stream__pair_8h_source.html | 142 ++ docs/gemm__traits_8h.html | 38 +- docs/gemm__traits_8h_source.html | 267 ++-- docs/globals.html | 11 +- docs/globals_defs.html | 5 +- docs/globals_func.html | 6 +- docs/group__IdentityBlockSwizzle.html | 84 + docs/group__fragment__concept.html | 5 +- docs/group__fragment__iterator__concept.html | 2 +- docs/group__layout__concept.html | 2 +- docs/group__predicate__iterator__concept.html | 2 +- docs/group__predicate__tile__adapter.html | 2 +- docs/group__predicate__vector__concept.html | 2 +- .../group__tile__load__iterator__concept.html | 4 +- ...group__tile__store__iterator__concept.html | 4 +- docs/group__tile__traits__concept.html | 4 +- docs/hgemm__global__tile_8h.html | 10 +- docs/hgemm__global__tile_8h_source.html | 16 +- docs/hgemm__multiply__add_8h.html | 10 +- docs/hgemm__multiply__add_8h_source.html | 31 +- docs/hgemm__swizzle_8h.html | 4 +- docs/hgemm__swizzle_8h_source.html | 8 +- docs/hgemm__traits_8h.html | 30 +- docs/hgemm__traits_8h_source.html | 122 +- docs/hierarchy.html | 685 +++++---- docs/igemm__epilogue_8h.html | 16 +- docs/igemm__epilogue_8h_source.html | 50 +- docs/igemm__global__tile_8h.html | 16 +- docs/igemm__global__tile_8h_source.html | 41 +- docs/igemm__multiply__add_8h.html | 10 +- docs/igemm__multiply__add_8h_source.html | 31 +- docs/igemm__swizzle_8h.html | 4 +- docs/igemm__swizzle_8h_source.html | 8 +- docs/igemm__traits_8h.html | 46 +- docs/igemm__traits_8h_source.html | 141 +- docs/index.html | 2 +- docs/iterator__access_8h.html | 70 +- docs/iterator__access_8h_source.html | 20 +- docs/kernel__launch_8h.html | 108 ++ docs/kernel__launch_8h_source.html | 95 ++ docs/linear__scaling_8h.html | 15 +- docs/linear__scaling_8h_source.html | 36 +- docs/linear__scaling__device__ptr_8h.html | 114 ++ ...inear__scaling__device__ptr_8h_source.html | 109 ++ docs/load__store_8h.html | 42 +- docs/load__store_8h_source.html | 83 +- docs/matrix__traits_8h.html | 28 +- docs/matrix__traits_8h_source.html | 91 +- docs/menudata.js | 17 + docs/modules.html | 17 +- docs/namespacecutlass.html | 1356 ++++++++--------- docs/namespacecutlass_1_1detail.html | 95 ++ docs/namespacecutlass_1_1gemm.html | 144 +- docs/namespacecutlass_1_1platform.html | 1185 +++++++++++++- docs/namespacemembers.html | 146 +- docs/namespacemembers_func.html | 144 +- docs/namespacemembers_type.html | 2 +- docs/namespaces.html | 7 +- docs/numeric__types_8h.html | 107 ++ docs/numeric__types_8h_source.html | 91 ++ docs/platform_8h.html | 4 +- docs/platform_8h_source.html | 166 +- docs/predicate__vector_8h.html | 11 +- docs/predicate__vector_8h_source.html | 128 +- docs/reshape__tile_8h.html | 4 +- docs/reshape__tile_8h_source.html | 6 +- docs/scalar__or__pointer_8h.html | 109 ++ docs/scalar__or__pointer_8h_source.html | 101 ++ docs/search/all_0.js | 1 + docs/search/all_1.js | 20 +- docs/search/all_10.js | 21 +- docs/search/all_11.js | 130 +- docs/search/all_12.js | 77 +- docs/search/all_13.js | 1 + docs/search/all_14.js | 17 +- docs/search/all_15.js | 5 +- docs/search/all_17.js | 9 +- docs/search/all_18.html | 30 + docs/search/all_18.js | 4 + docs/search/all_2.js | 19 +- docs/search/all_3.js | 55 +- docs/search/all_4.js | 21 +- docs/search/all_5.js | 7 +- docs/search/all_6.js | 45 +- docs/search/all_7.js | 87 +- docs/search/all_8.js | 5 +- docs/search/all_9.js | 71 +- docs/search/all_a.js | 93 +- docs/search/all_b.js | 40 +- docs/search/all_c.js | 27 +- docs/search/all_d.js | 7 +- docs/search/all_e.js | 53 +- docs/search/all_f.js | 22 +- docs/search/classes_1.js | 1 + docs/search/classes_10.js | 33 +- docs/search/classes_11.js | 9 +- docs/search/classes_12.js | 15 +- docs/search/classes_13.html | 30 + docs/search/classes_13.js | 5 + docs/search/classes_14.html | 30 + docs/search/classes_14.js | 8 + docs/search/classes_2.js | 18 +- docs/search/classes_3.js | 5 +- docs/search/classes_5.js | 11 +- docs/search/classes_6.js | 20 +- docs/search/classes_8.js | 17 +- docs/search/classes_9.js | 12 +- docs/search/classes_a.js | 22 +- docs/search/classes_b.js | 8 +- docs/search/classes_c.js | 13 +- docs/search/classes_d.js | 23 +- docs/search/classes_e.js | 45 +- docs/search/classes_f.js | 62 +- docs/search/defines_1.js | 1 + docs/search/enums_0.js | 2 +- docs/search/enumvalues_1.js | 17 +- docs/search/enumvalues_2.js | 19 +- docs/search/enumvalues_3.js | 2 +- docs/search/enumvalues_4.html | 30 + docs/search/enumvalues_4.js | 4 + docs/search/enumvalues_5.html | 30 + docs/search/enumvalues_5.js | 4 + docs/search/files_0.js | 1 + docs/search/files_10.html | 30 + docs/search/files_10.js | 6 + docs/search/files_2.js | 3 +- docs/search/files_3.js | 4 + docs/search/files_5.js | 1 - docs/search/files_6.js | 3 +- docs/search/files_7.js | 4 +- docs/search/files_8.js | 3 +- docs/search/files_9.js | 2 +- docs/search/files_a.js | 4 +- docs/search/files_b.js | 6 +- docs/search/files_c.js | 4 +- docs/search/files_d.js | 16 +- docs/search/files_e.js | 7 +- docs/search/files_f.html | 30 + docs/search/files_f.js | 9 + docs/search/functions_0.js | 3 +- docs/search/functions_1.js | 7 +- docs/search/functions_10.js | 19 +- docs/search/functions_11.js | 21 +- docs/search/functions_12.js | 14 +- docs/search/functions_13.js | 2 +- docs/search/functions_14.js | 3 +- docs/search/functions_15.html | 30 + docs/search/functions_15.js | 5 + docs/search/functions_16.html | 30 + docs/search/functions_16.js | 7 + docs/search/functions_17.html | 30 + docs/search/functions_17.js | 4 + docs/search/functions_2.js | 4 +- docs/search/functions_3.js | 31 +- docs/search/functions_4.js | 8 +- docs/search/functions_5.js | 7 +- docs/search/functions_6.js | 12 +- docs/search/functions_7.js | 29 +- docs/search/functions_8.js | 5 +- docs/search/functions_9.js | 31 +- docs/search/functions_a.js | 11 +- docs/search/functions_b.js | 17 +- docs/search/functions_c.js | 43 +- docs/search/functions_d.js | 6 +- docs/search/functions_e.js | 38 +- docs/search/functions_f.js | 22 +- docs/search/groups_1.js | 2 +- docs/search/groups_2.js | 4 +- docs/search/groups_3.js | 6 +- docs/search/groups_4.html | 30 + docs/search/groups_4.js | 6 + docs/search/namespaces_0.js | 1 + docs/search/searchdata.js | 16 +- docs/search/typedefs_0.js | 9 +- docs/search/typedefs_1.js | 7 +- docs/search/typedefs_10.js | 4 +- docs/search/typedefs_11.js | 2 +- docs/search/typedefs_12.html | 30 + docs/search/typedefs_12.js | 4 + docs/search/typedefs_2.js | 9 +- docs/search/typedefs_3.js | 3 +- docs/search/typedefs_4.js | 2 +- docs/search/typedefs_5.js | 19 +- docs/search/typedefs_6.js | 32 +- docs/search/typedefs_7.js | 14 +- docs/search/typedefs_8.js | 2 +- docs/search/typedefs_9.js | 5 +- docs/search/typedefs_a.js | 5 +- docs/search/typedefs_b.js | 4 +- docs/search/typedefs_c.js | 6 +- docs/search/typedefs_d.js | 32 +- docs/search/typedefs_e.js | 52 +- docs/search/typedefs_f.js | 29 +- docs/search/variables_0.js | 5 +- docs/search/variables_1.js | 9 +- docs/search/variables_10.js | 9 +- docs/search/variables_11.html | 30 + docs/search/variables_11.js | 5 + docs/search/variables_2.js | 1 + docs/search/variables_3.js | 6 +- docs/search/variables_5.js | 8 +- docs/search/variables_6.js | 5 +- docs/search/variables_7.js | 16 +- docs/search/variables_8.js | 74 +- docs/search/variables_9.js | 8 +- docs/search/variables_a.js | 5 +- docs/search/variables_b.js | 2 +- docs/search/variables_c.js | 7 +- docs/search/variables_d.js | 10 +- docs/search/variables_e.js | 18 +- docs/search/variables_f.js | 22 +- docs/sgemm__traits_8h.html | 23 +- docs/sgemm__traits_8h_source.html | 19 +- docs/shape_8h.html | 20 +- docs/shape_8h_source.html | 23 +- docs/structDebugType.html | 88 ++ docs/structDebugValue.html | 88 ++ docs/structcutlass_1_1AlignedStruct.html | 2 +- ...ass_1_1ComputeOffsetFromShape-members.html | 4 +- ...ructcutlass_1_1ComputeOffsetFromShape.html | 12 +- ...s_1_1ComputeOffsetFromStrides-members.html | 4 +- ...ctcutlass_1_1ComputeOffsetFromStrides.html | 12 +- ...omputeThreadOffsetFromStrides-members.html | 2 +- ...ass_1_1ComputeThreadOffsetFromStrides.html | 2 +- ...011_03ed682791cf043da79a7cc93228a8c85.html | 2 +- ...011_0e75281d7e02fa191f5d498e10e25dc1b.html | 2 +- ...01T__dd54c41f6edb97d3c208cb7c6fe4ab9b.html | 2 +- ...01T__f2e6d84a53db391977c787a65ed62aca.html | 2 +- ..._1_1ConstPredicateTileAdapter-members.html | 2 +- ...tcutlass_1_1ConstPredicateTileAdapter.html | 2 +- docs/structcutlass_1_1Convert.html | 2 +- ...3_01Ob568b5e19b6f78a5fa50d1f821f0bc2a.html | 2 +- ...3_01Ofca5985d18bcb54bc1f49355f3cee121.html | 2 +- docs/structcutlass_1_1Coord-members.html | 62 +- docs/structcutlass_1_1Coord.html | 633 +++++--- docs/structcutlass_1_1Copy-members.html | 2 +- docs/structcutlass_1_1Copy.html | 2 +- docs/structcutlass_1_1DumpType.html | 92 ++ docs/structcutlass_1_1Extent-members.html | 2 +- docs/structcutlass_1_1Extent.html | 2 +- ...or_3_01T_00_01Lanes_01_4_01_4-members.html | 2 +- ...3_01Vector_3_01T_00_01Lanes_01_4_01_4.html | 2 +- ..._00_01Lanes_01_4_01const_01_4-members.html | 2 +- ...or_3_01T_00_01Lanes_01_4_01const_01_4.html | 2 +- docs/structcutlass_1_1Fragment-members.html | 13 +- docs/structcutlass_1_1Fragment.html | 77 +- ...lass_1_1FragmentConstIterator-members.html | 14 +- ...tructcutlass_1_1FragmentConstIterator.html | 76 +- ...utlass_1_1FragmentElementType-members.html | 93 ++ .../structcutlass_1_1FragmentElementType.html | 125 ++ ...ctcutlass_1_1FragmentIterator-members.html | 16 +- docs/structcutlass_1_1FragmentIterator.html | 88 +- .../structcutlass_1_1GemmOperand-members.html | 2 +- docs/structcutlass_1_1GemmOperand.html | 2 +- docs/structcutlass_1_1Identity-members.html | 2 +- docs/structcutlass_1_1Identity.html | 2 +- ...lass_1_1IdentityTensorMapFunc-members.html | 92 ++ ...tructcutlass_1_1IdentityTensorMapFunc.html | 166 ++ ...uctcutlass_1_1IteratorAdvance-members.html | 2 +- docs/structcutlass_1_1IteratorAdvance.html | 2 +- ..._1_1KernelLaunchConfiguration-members.html | 94 ++ ...tcutlass_1_1KernelLaunchConfiguration.html | 213 +++ ...ctcutlass_1_1KernelLaunchConfiguration.png | Bin 0 -> 2229 bytes docs/structcutlass_1_1Load-members.html | 10 +- docs/structcutlass_1_1Load.html | 36 +- ...1_1kS1b28106546bd22002cd52e60197548ce.html | 171 +++ ...1_1kS3de71ba9ef724a37fb1cf315da8bf9b5.html | 171 +++ ...1_1kS840de374cd1e3ee3a4aa6a70a9fa83d1.html | 92 ++ ...1_1kS91ff65dd36f118bca4542df6128d6d99.html | 171 +++ ...1_1kS966cdf9ba449f056a92458bf87878053.html | 92 ++ ...1_1kSc8e65d97c8f615e98ac2e03f0bcd236e.html | 92 ++ ...1_1kSf1f63874ddc1302ed9e60e0478e1a8ad.html | 174 +++ ...1_1kSf2952bcb62d20c76c595eac4c59f7239.html | 92 ++ ...1_1kW847237836867f08e1121b00b7e44d8ae.html | 92 ++ ...1_1kWe58d0048e24352beeec002fd483c53b2.html | 171 +++ ..._01Fr16ad55dd7bca84e6a7fc608c26eee889.html | 92 ++ ..._01Fra240d9c28383cd3945277ec3a927c538.html | 171 +++ ...01Fra03475418f27732d726d18feb23feeed2.html | 92 ++ ...01Fra934529165fa8fecdd392b5302d25ef26.html | 171 +++ ..._01Fr23088b868c6eeec9377c46892553686f.html | 171 +++ ..._01Frc8b0aa6ddd9b4317158c26574a6881de.html | 92 ++ ...0_01d9e675253ca19588f1ae4bd898579523f.html | 92 ++ ...00_01double_00_01kStride_00_0116_01_4.html | 171 +++ .../structcutlass_1_1MatrixCoord-members.html | 140 ++ docs/structcutlass_1_1MatrixCoord.html | 748 +++++++++ docs/structcutlass_1_1MatrixCoord.png | Bin 0 -> 640 bytes ...structcutlass_1_1MatrixLayout-members.html | 2 +- docs/structcutlass_1_1MatrixLayout.html | 25 +- ..._1MatrixLayout_1_1ColumnMajor-members.html | 92 ++ ...utlass_1_1MatrixLayout_1_1ColumnMajor.html | 161 ++ ...out_1_1ColumnMajorBlockLinear-members.html | 95 ++ ...atrixLayout_1_1ColumnMajorBlockLinear.html | 257 ++++ ...out_1_1ColumnMajorInterleaved-members.html | 94 ++ ...atrixLayout_1_1ColumnMajorInterleaved.html | 230 +++ ...rixLayout_1_1ContiguousLayout-members.html | 95 ++ ...s_1_1MatrixLayout_1_1ContiguousLayout.html | 254 +++ ...s_1_1MatrixLayout_1_1RowMajor-members.html | 92 ++ ...ctcutlass_1_1MatrixLayout_1_1RowMajor.html | 161 ++ ...Layout_1_1RowMajorBlockLinear-members.html | 95 ++ ..._1MatrixLayout_1_1RowMajorBlockLinear.html | 257 ++++ ...Layout_1_1RowMajorInterleaved-members.html | 94 ++ ..._1MatrixLayout_1_1RowMajorInterleaved.html | 230 +++ ...uctcutlass_1_1MatrixTransform-members.html | 93 ++ docs/structcutlass_1_1MatrixTransform.html | 126 ++ docs/structcutlass_1_1Max-members.html | 91 ++ docs/structcutlass_1_1Max.html | 126 ++ .../structcutlass_1_1MemorySpace-members.html | 2 +- docs/structcutlass_1_1MemorySpace.html | 2 +- docs/structcutlass_1_1Min-members.html | 91 ++ docs/structcutlass_1_1Min.html | 126 ++ ...tlass_1_1PredicateTileAdapter-members.html | 2 +- ...structcutlass_1_1PredicateTileAdapter.html | 2 +- ...uctcutlass_1_1PredicateVector-members.html | 2 +- docs/structcutlass_1_1PredicateVector.html | 2 +- ...cateVector_1_1TrivialIterator-members.html | 2 +- ...1_1PredicateVector_1_1TrivialIterator.html | 2 +- ...s_1_1PredicatedTileLoadStream-members.html | 111 ++ ...ctcutlass_1_1PredicatedTileLoadStream.html | 402 +++++ ...uctcutlass_1_1PredicatedTileLoadStream.png | Bin 0 -> 1349 bytes ..._1_1PredicatedTileStoreStream-members.html | 112 ++ ...tcutlass_1_1PredicatedTileStoreStream.html | 466 ++++++ ...ctcutlass_1_1PredicatedTileStoreStream.png | Bin 0 -> 1350 bytes ..._1RegularTilePredicateFunctor-members.html | 94 ++ ...utlass_1_1RegularTilePredicateFunctor.html | 223 +++ .../structcutlass_1_1ReshapeTile-members.html | 2 +- docs/structcutlass_1_1ReshapeTile.html | 2 +- ...1kAccessSize___00_01true_01_4-members.html | 2 +- ...e___00_01kAccessSize___00_01true_01_4.html | 2 +- docs/structcutlass_1_1ScalarIO-members.html | 93 ++ docs/structcutlass_1_1ScalarIO.html | 188 +++ docs/structcutlass_1_1Shape-members.html | 2 +- docs/structcutlass_1_1Shape.html | 2 +- docs/structcutlass_1_1ShapeAdd-members.html | 2 +- docs/structcutlass_1_1ShapeAdd.html | 2 +- docs/structcutlass_1_1ShapeCount-members.html | 2 +- docs/structcutlass_1_1ShapeCount.html | 2 +- docs/structcutlass_1_1ShapeDiv-members.html | 2 +- docs/structcutlass_1_1ShapeDiv.html | 2 +- ...uctcutlass_1_1ShapeDivCeiling-members.html | 91 ++ docs/structcutlass_1_1ShapeDivCeiling.html | 118 ++ docs/structcutlass_1_1ShapeMax-members.html | 2 +- docs/structcutlass_1_1ShapeMax.html | 2 +- docs/structcutlass_1_1ShapeMin-members.html | 2 +- docs/structcutlass_1_1ShapeMin.html | 2 +- docs/structcutlass_1_1ShapeMul-members.html | 2 +- docs/structcutlass_1_1ShapeMul.html | 2 +- docs/structcutlass_1_1ShapeScale-members.html | 2 +- docs/structcutlass_1_1ShapeScale.html | 2 +- ...structcutlass_1_1ShapeStrides-members.html | 8 +- docs/structcutlass_1_1ShapeStrides.html | 18 +- docs/structcutlass_1_1ShapeSub-members.html | 2 +- docs/structcutlass_1_1ShapeSub.html | 2 +- .../structcutlass_1_1StorageType-members.html | 8 +- docs/structcutlass_1_1StorageType.html | 18 +- ...ass_1_1StorageType_3_011_01_4-members.html | 2 +- ...ructcutlass_1_1StorageType_3_011_01_4.html | 2 +- ...ass_1_1StorageType_3_012_01_4-members.html | 2 +- ...ructcutlass_1_1StorageType_3_012_01_4.html | 2 +- ...ass_1_1StorageType_3_014_01_4-members.html | 2 +- ...ructcutlass_1_1StorageType_3_014_01_4.html | 2 +- docs/structcutlass_1_1Store-members.html | 10 +- docs/structcutlass_1_1Store.html | 36 +- ..._1_1k004b304998a534d76357f834068909f8.html | 92 ++ ..._1_1k12f5c8a016a307e76de374322fc00a66.html | 92 ++ ..._1_1k220d5790f803f10840e2a92fb9a51dac.html | 171 +++ ..._1_1k28cc0b88a16efca73d258128312d2a7e.html | 171 +++ ..._1_1k40d038d4bce377843c21a56ebf97d011.html | 171 +++ ..._1_1k60eedca420c41e94fd40b41299967ef2.html | 92 ++ ..._1_1k775a1d27affec5236489735ed4503c92.html | 92 ++ ..._1_1k84da7dcd68ee74b8d2bdb67885b0ca56.html | 92 ++ ..._1_1kd9a7e85f80a21c504388612a60462417.html | 171 +++ ..._1_1ke6d73d34fa7b5254cf828804a19842e1.html | 171 +++ ...00_013d38935f41bf709e067932b9e042255a.html | 92 ++ ...00_0160391c6be5cb1d3f99e012a6a18e486d.html | 171 +++ ...ructcutlass_1_1TensorRefArray-members.html | 101 ++ docs/structcutlass_1_1TensorRefArray.html | 398 +++++ ...lass_1_1TensorRefBatchStrided-members.html | 134 ++ ...tructcutlass_1_1TensorRefBatchStrided.html | 537 +++++++ ...structcutlass_1_1TensorRefBatchStrided.png | Bin 0 -> 1810 bytes ...00_01LongIndex___01_4_1_1StrideVector.html | 97 ++ ...ructcutlass_1_1TileAllocation-members.html | 101 ++ docs/structcutlass_1_1TileAllocation.html | 366 +++++ docs/structcutlass_1_1TileCoord-members.html | 151 ++ docs/structcutlass_1_1TileCoord.html | 1127 ++++++++++++++ docs/structcutlass_1_1TileCoord.png | Bin 0 -> 720 bytes ...ctcutlass_1_1TileIteratorBase-members.html | 54 +- docs/structcutlass_1_1TileIteratorBase.html | 402 ++--- docs/structcutlass_1_1TileIteratorBase.png | Bin 2958 -> 2987 bytes ...1_1TileIteratorBase_1_1Params-members.html | 30 +- ...cutlass_1_1TileIteratorBase_1_1Params.html | 328 +++- ...tcutlass_1_1TileIteratorBase_1_1Params.png | Bin 4162 -> 4455 bytes ...ctcutlass_1_1TileLoadIterator-members.html | 102 +- docs/structcutlass_1_1TileLoadIterator.html | 1015 +++++++----- docs/structcutlass_1_1TileLoadIterator.png | Bin 2304 -> 2306 bytes ...1_1TileLoadIterator_1_1Params-members.html | 44 +- ...cutlass_1_1TileLoadIterator_1_1Params.html | 499 ++++-- ...tcutlass_1_1TileLoadIterator_1_1Params.png | Bin 3293 -> 3444 bytes ...ructcutlass_1_1TileLoadStream-members.html | 106 ++ docs/structcutlass_1_1TileLoadStream.html | 525 +++++++ docs/structcutlass_1_1TileLoadStream.png | Bin 0 -> 1352 bytes ...s_1_1TileLoadStream_1_1Params-members.html | 93 ++ ...ctcutlass_1_1TileLoadStream_1_1Params.html | 188 +++ ..._1_1TileLoadStream_1_1PredicateVector.html | 95 ++ ...tcutlass_1_1TileStoreIterator-members.html | 103 +- docs/structcutlass_1_1TileStoreIterator.html | 1205 ++++++++++----- docs/structcutlass_1_1TileStoreIterator.png | Bin 2309 -> 2318 bytes ..._1TileStoreIterator_1_1Params-members.html | 43 +- ...utlass_1_1TileStoreIterator_1_1Params.html | 424 +++++- ...cutlass_1_1TileStoreIterator_1_1Params.png | Bin 2388 -> 2515 bytes ...uctcutlass_1_1TileStoreStream-members.html | 107 ++ docs/structcutlass_1_1TileStoreStream.html | 556 +++++++ docs/structcutlass_1_1TileStoreStream.png | Bin 0 -> 1351 bytes ..._1_1TileStoreStream_1_1Params-members.html | 93 ++ ...tcutlass_1_1TileStoreStream_1_1Params.html | 188 +++ ...1_1TileStoreStream_1_1PredicateVector.html | 95 ++ docs/structcutlass_1_1TileTraits-members.html | 16 +- docs/structcutlass_1_1TileTraits.html | 113 +- ..._1_1TileTraitsContiguousMajor-members.html | 2 +- ...tcutlass_1_1TileTraitsContiguousMajor.html | 2 +- ...cutlass_1_1TileTraitsStandard-members.html | 11 +- docs/structcutlass_1_1TileTraitsStandard.html | 29 +- ...lass_1_1TileTraitsStrideMajor-members.html | 2 +- ...tructcutlass_1_1TileTraitsStrideMajor.html | 2 +- ...cutlass_1_1TileTraitsWarpRake-members.html | 2 +- docs/structcutlass_1_1TileTraitsWarpRake.html | 2 +- ...raitsWarpRake_1_1ThreadOffset-members.html | 2 +- ...1_1TileTraitsWarpRake_1_1ThreadOffset.html | 2 +- ...tcutlass_1_1TiledThreadOffset-members.html | 2 +- docs/structcutlass_1_1TiledThreadOffset.html | 2 +- ..._1TrivialPredicateTileAdapter-members.html | 2 +- ...utlass_1_1TrivialPredicateTileAdapter.html | 2 +- ...structcutlass_1_1VectorTraits-members.html | 2 +- docs/structcutlass_1_1VectorTraits.html | 2 +- ...or_3_01T_00_01Lanes_01_4_01_4-members.html | 2 +- ...3_01Vector_3_01T_00_01Lanes_01_4_01_4.html | 2 +- ..._00_01Lanes_01_4_01const_01_4-members.html | 2 +- ...or_3_01T_00_01Lanes_01_4_01const_01_4.html | 2 +- docs/structcutlass_1_1Vectorize-members.html | 2 +- docs/structcutlass_1_1Vectorize.html | 2 +- ..._0132_01_4_00_01kLanes___01_4-members.html | 91 ++ ...n1__t_00_0132_01_4_00_01kLanes___01_4.html | 118 ++ ...0_018_01_4_00_01kLanes___01_4-members.html | 91 ++ ...nt4__t_00_018_01_4_00_01kLanes___01_4.html | 118 ++ ...0_018_01_4_00_01kLanes___01_4-members.html | 91 ++ ...nt4__t_00_018_01_4_00_01kLanes___01_4.html | 118 ++ docs/structcutlass_1_1ZipConvert-members.html | 99 ++ docs/structcutlass_1_1ZipConvert.html | 340 +++++ .../structcutlass_1_1ZipFragment-members.html | 98 ++ docs/structcutlass_1_1ZipFragment.html | 310 ++++ ...structcutlass_1_1ZipTensorRef-members.html | 96 ++ docs/structcutlass_1_1ZipTensorRef.html | 255 ++++ ...tcutlass_1_1ZipTileAllocation-members.html | 98 ++ docs/structcutlass_1_1ZipTileAllocation.html | 287 ++++ ..._1_1ZipTileIterator_1_1Params-members.html | 94 ++ ...tcutlass_1_1ZipTileIterator_1_1Params.html | 217 +++ docs/structcutlass_1_1bin1__t.html | 92 ++ ...ructcutlass_1_1divide__assert-members.html | 2 +- docs/structcutlass_1_1divide__assert.html | 2 +- ..._1_1gemm_1_1ClearAccumulators-members.html | 3 +- ...tcutlass_1_1gemm_1_1ClearAccumulators.html | 34 +- ...1_1ClearAccumulators_1_1SharedStorage.html | 2 +- ...mm_1_1ColumnMajorBlockSwizzle-members.html | 95 ++ ...ss_1_1gemm_1_1ColumnMajorBlockSwizzle.html | 260 ++++ ...utlass_1_1gemm_1_1DgemmConfig-members.html | 59 +- .../structcutlass_1_1gemm_1_1DgemmConfig.html | 163 +- docs/structcutlass_1_1gemm_1_1DgemmConfig.png | Bin 2957 -> 3040 bytes ...utlass_1_1gemm_1_1DgemmTraits-members.html | 19 +- .../structcutlass_1_1gemm_1_1DgemmTraits.html | 37 +- docs/structcutlass_1_1gemm_1_1DgemmTraits.png | Bin 8483 -> 8474 bytes ...ss_1_1gemm_1_1Fp16SgemmConfig-members.html | 118 ++ ...uctcutlass_1_1gemm_1_1Fp16SgemmConfig.html | 186 +++ ...ructcutlass_1_1gemm_1_1Fp16SgemmConfig.png | Bin 0 -> 3199 bytes ...1gemm_1_1Fp16SgemmSgemmTraits-members.html | 114 ++ ...tlass_1_1gemm_1_1Fp16SgemmSgemmTraits.html | 182 +++ ...utlass_1_1gemm_1_1Fp16SgemmSgemmTraits.png | Bin 0 -> 8749 bytes ..._1gemm_1_1FragmentMultiplyAdd-members.html | 19 +- ...utlass_1_1gemm_1_1FragmentMultiplyAdd.html | 131 +- ...half_00_01half_00_01true_01_4-members.html | 96 ++ ...Add_3_01half_00_01half_00_01true_01_4.html | 287 ++++ ...structcutlass_1_1gemm_1_1Gemm-members.html | 12 +- docs/structcutlass_1_1gemm_1_1Gemm.html | 151 +- ...cutlass_1_1gemm_1_1GemmConfig-members.html | 59 +- docs/structcutlass_1_1gemm_1_1GemmConfig.html | 427 +++--- ...tcutlass_1_1gemm_1_1GemmCoord-members.html | 152 ++ docs/structcutlass_1_1gemm_1_1GemmCoord.html | 1102 ++++++++++++++ docs/structcutlass_1_1gemm_1_1GemmCoord.png | Bin 0 -> 691 bytes ...ctcutlass_1_1gemm_1_1GemmDesc-members.html | 44 +- docs/structcutlass_1_1gemm_1_1GemmDesc.html | 669 ++++++-- ...tlass_1_1gemm_1_1GemmEpilogue-members.html | 41 +- ...structcutlass_1_1gemm_1_1GemmEpilogue.html | 147 +- ...1_1gemm_1_1GemmEpilogueTraits-members.html | 38 +- ...cutlass_1_1gemm_1_1GemmEpilogueTraits.html | 228 +-- ...m_1_1GemmEpilogueTraitsHelper-members.html | 11 +- ...s_1_1gemm_1_1GemmEpilogueTraitsHelper.html | 65 +- ...1GemmEpilogueTraits_1_1Params-members.html | 22 +- ...1gemm_1_1GemmEpilogueTraits_1_1Params.html | 118 +- ...ilogueTraits_1_1SharedStorage-members.html | 9 +- ..._1GemmEpilogueTraits_1_1SharedStorage.html | 54 +- ...1gemm_1_1GemmGlobalIteratorAb-members.html | 109 +- ...tlass_1_1gemm_1_1GemmGlobalIteratorAb.html | 723 ++++++--- ...utlass_1_1gemm_1_1GemmGlobalIteratorAb.png | Bin 3492 -> 4486 bytes ...emmGlobalIteratorAb_1_1Params-members.html | 42 +- ...emm_1_1GemmGlobalIteratorAb_1_1Params.html | 135 +- ...gemm_1_1GemmGlobalIteratorAb_1_1Params.png | Bin 3308 -> 3459 bytes ...1gemm_1_1GemmGlobalIteratorCd-members.html | 72 +- ...tlass_1_1gemm_1_1GemmGlobalIteratorCd.html | 531 +++++-- ...utlass_1_1gemm_1_1GemmGlobalIteratorCd.png | Bin 1734 -> 2546 bytes ...emmGlobalIteratorCd_1_1Params-members.html | 7 +- ...emm_1_1GemmGlobalIteratorCd_1_1Params.html | 48 +- ...gemm_1_1GemmGlobalIteratorCd_1_1Params.png | Bin 0 -> 1464 bytes ...emm_1_1GemmGlobalTileCdTraits-members.html | 5 +- ...ass_1_1gemm_1_1GemmGlobalTileCdTraits.html | 43 +- ...lTileCdTraits_1_1ThreadOffset-members.html | 2 +- ...emmGlobalTileCdTraits_1_1ThreadOffset.html | 2 +- ...1gemm_1_1GemmGlobalTileTraits-members.html | 11 +- ...tlass_1_1gemm_1_1GemmGlobalTileTraits.html | 75 +- ...utlass_1_1gemm_1_1GemmGlobalTileTraits.png | Bin 2856 -> 2833 bytes ...balTileTraits_1_1ThreadOffset-members.html | 2 +- ...1GemmGlobalTileTraits_1_1ThreadOffset.html | 2 +- ...emm_1_1GemmMultiplicandTraits-members.html | 2 +- ...ass_1_1gemm_1_1GemmMultiplicandTraits.html | 2 +- ..._1gemm_1_1GemmOperandTraitsAb-members.html | 2 +- ...utlass_1_1gemm_1_1GemmOperandTraitsAb.html | 2 +- ..._1_1GemmSharedLoadTileATraits-members.html | 6 +- ..._1_1gemm_1_1GemmSharedLoadTileATraits.html | 25 +- ...adTileATraits_1_1ThreadOffset-members.html | 2 +- ...SharedLoadTileATraits_1_1ThreadOffset.html | 2 +- ..._1_1GemmSharedLoadTileBTraits-members.html | 6 +- ..._1_1gemm_1_1GemmSharedLoadTileBTraits.html | 25 +- ...adTileBTraits_1_1ThreadOffset-members.html | 2 +- ...SharedLoadTileBTraits_1_1ThreadOffset.html | 2 +- ..._1_1GemmSharedLoadTileDTraits-members.html | 23 +- ..._1_1gemm_1_1GemmSharedLoadTileDTraits.html | 66 +- ...adTileDTraits_1_1ThreadOffset-members.html | 2 +- ...SharedLoadTileDTraits_1_1ThreadOffset.html | 2 +- ..._1GemmSharedStoreTileAbTraits-members.html | 2 +- ..._1gemm_1_1GemmSharedStoreTileAbTraits.html | 2 +- ...eTileAbTraits_1_1ThreadOffset-members.html | 2 +- ...aredStoreTileAbTraits_1_1ThreadOffset.html | 2 +- ...1_1GemmSharedStoreTileDTraits-members.html | 2 +- ...1_1gemm_1_1GemmSharedStoreTileDTraits.html | 2 +- ...reTileDTraits_1_1ThreadOffset-members.html | 2 +- ...haredStoreTileDTraits_1_1ThreadOffset.html | 2 +- ...aredStoreWithSkewTileAbTraits-members.html | 2 +- ..._1GemmSharedStoreWithSkewTileAbTraits.html | 2 +- ...wTileAbTraits_1_1ThreadOffset-members.html | 2 +- ...eWithSkewTileAbTraits_1_1ThreadOffset.html | 2 +- ...lass_1_1gemm_1_1GemmTileTraitsHelperA.html | 2 +- ...nMajor_00_01GemmConfig___01_4-members.html | 2 +- ..._1kColumnMajor_00_01GemmConfig___01_4.html | 8 +- ...1_1kColumnMajor_00_01GemmConfig___01_4.png | Bin 1679 -> 1846 bytes ...wMajor_00_01GemmConfig___01_4-members.html | 9 +- ...t_1_1kRowMajor_00_01GemmConfig___01_4.html | 67 +- ...lass_1_1gemm_1_1GemmTileTraitsHelperB.html | 2 +- ...nMajor_00_01GemmConfig___01_4-members.html | 9 +- ..._1kColumnMajor_00_01GemmConfig___01_4.html | 67 +- ...wMajor_00_01GemmConfig___01_4-members.html | 2 +- ...t_1_1kRowMajor_00_01GemmConfig___01_4.html | 8 +- ...ut_1_1kRowMajor_00_01GemmConfig___01_4.png | Bin 1637 -> 1731 bytes ...cutlass_1_1gemm_1_1GemmTraits-members.html | 15 +- docs/structcutlass_1_1gemm_1_1GemmTraits.html | 143 +- ...aits_1_1MainLoopSharedStorage-members.html | 6 +- ..._1GemmTraits_1_1MainLoopSharedStorage.html | 31 +- ...1gemm_1_1GemmTraits_1_1Params-members.html | 18 +- ...tlass_1_1gemm_1_1GemmTraits_1_1Params.html | 393 +++-- ...utlass_1_1gemm_1_1GemmTraits_1_1Params.png | Bin 0 -> 2212 bytes docs/structcutlass_1_1gemm_1_1GetExtent.html | 2 +- ...perand_1_1kA_00_01Tile___01_4-members.html | 2 +- ..._01GemmOperand_1_1kA_00_01Tile___01_4.html | 2 +- ...perand_1_1kB_00_01Tile___01_4-members.html | 2 +- ..._01GemmOperand_1_1kB_00_01Tile___01_4.html | 2 +- ...s_1_1gemm_1_1GlobalLoadStream-members.html | 58 +- ...ctcutlass_1_1gemm_1_1GlobalLoadStream.html | 809 ++++++++-- ...1gemm_1_1GlobalLoadStreamPair-members.html | 104 ++ ...tlass_1_1gemm_1_1GlobalLoadStreamPair.html | 519 +++++++ ...lobalLoadStreamPair_1_1Params-members.html | 94 ++ ...emm_1_1GlobalLoadStreamPair_1_1Params.html | 217 +++ ...adStreamPair_1_1SharedStorage-members.html | 92 ++ ...GlobalLoadStreamPair_1_1SharedStorage.html | 139 ++ ...1_1GlobalLoadStream_1_1Params-members.html | 94 ++ ...1_1gemm_1_1GlobalLoadStream_1_1Params.html | 215 +++ ..._1_1GlobalLoadStream_1_1SharedStorage.html | 97 ++ ...utlass_1_1gemm_1_1HgemmConfig-members.html | 59 +- .../structcutlass_1_1gemm_1_1HgemmConfig.html | 163 +- docs/structcutlass_1_1gemm_1_1HgemmConfig.png | Bin 2911 -> 2958 bytes ...gemmCrosswiseGlobalTileTraits-members.html | 9 +- ...emm_1_1HgemmCrosswiseGlobalTileTraits.html | 59 +- ...balTileTraits_1_1ThreadOffset-members.html | 2 +- ...swiseGlobalTileTraits_1_1ThreadOffset.html | 2 +- ...tlass_1_1gemm_1_1HgemmSwizzle-members.html | 2 +- ...structcutlass_1_1gemm_1_1HgemmSwizzle.html | 2 +- ...ass_1_1gemm_1_1HgemmTileTraitsHelperA.html | 2 +- ...wMajor_00_01GemmConfig___01_4-members.html | 9 +- ...t_1_1kRowMajor_00_01GemmConfig___01_4.html | 55 +- ...ass_1_1gemm_1_1HgemmTileTraitsHelperB.html | 2 +- ...nMajor_00_01GemmConfig___01_4-members.html | 9 +- ..._1kColumnMajor_00_01GemmConfig___01_4.html | 55 +- ...utlass_1_1gemm_1_1HgemmTraits-members.html | 19 +- .../structcutlass_1_1gemm_1_1HgemmTraits.html | 37 +- docs/structcutlass_1_1gemm_1_1HgemmTraits.png | Bin 3190 -> 3169 bytes ..._1_1gemm_1_1HgemmTraitsHelper-members.html | 44 +- ...tcutlass_1_1gemm_1_1HgemmTraitsHelper.html | 270 ++-- ...tcutlass_1_1gemm_1_1HgemmTransformerA.html | 2 +- ...umnMajor_00_01Iterator___01_4-members.html | 2 +- ..._1_1kColumnMajor_00_01Iterator___01_4.html | 2 +- ...RowMajor_00_01Iterator___01_4-members.html | 2 +- ...out_1_1kRowMajor_00_01Iterator___01_4.html | 2 +- ...tcutlass_1_1gemm_1_1HgemmTransformerB.html | 2 +- ...umnMajor_00_01Iterator___01_4-members.html | 2 +- ..._1_1kColumnMajor_00_01Iterator___01_4.html | 2 +- ...RowMajor_00_01Iterator___01_4-members.html | 2 +- ...out_1_1kRowMajor_00_01Iterator___01_4.html | 2 +- ...1gemm_1_1IdentityBlockSwizzle-members.html | 9 +- ...tlass_1_1gemm_1_1IdentityBlockSwizzle.html | 111 +- ...utlass_1_1gemm_1_1IgemmConfig-members.html | 59 +- .../structcutlass_1_1gemm_1_1IgemmConfig.html | 163 +- docs/structcutlass_1_1gemm_1_1IgemmConfig.png | Bin 2603 -> 2581 bytes ...t_00_01ThreadGemmShape___01_4-members.html | 118 ++ ...01int8__t_00_01ThreadGemmShape___01_4.html | 186 +++ ..._01int8__t_00_01ThreadGemmShape___01_4.png | Bin 0 -> 2442 bytes ...lass_1_1gemm_1_1IgemmEpilogue-members.html | 43 +- ...tructcutlass_1_1gemm_1_1IgemmEpilogue.html | 59 +- ..._1gemm_1_1IgemmEpilogueScalar-members.html | 2 +- ...utlass_1_1gemm_1_1IgemmEpilogueScalar.html | 2 +- ...mmEpilogueScalar_3_01int_01_4-members.html | 2 +- ...m_1_1IgemmEpilogueScalar_3_01int_01_4.html | 2 +- ..._1gemm_1_1IgemmEpilogueTraits-members.html | 34 +- ...utlass_1_1gemm_1_1IgemmEpilogueTraits.html | 100 +- ...cutlass_1_1gemm_1_1IgemmEpilogueTraits.png | Bin 3581 -> 3632 bytes ..._1_1IgemmEpilogueTraitsHelper-members.html | 13 +- ..._1_1gemm_1_1IgemmEpilogueTraitsHelper.html | 47 +- ...ilogueTraits___00_01true_01_4-members.html | 43 +- ...01GemmEpilogueTraits___00_01true_01_4.html | 59 +- ..._1_1IgemmFloatToInt8Converter-members.html | 2 +- ..._1_1gemm_1_1IgemmFloatToInt8Converter.html | 2 +- ...gemm_1_1IgemmGlobalIteratorAb-members.html | 151 ++ ...lass_1_1gemm_1_1IgemmGlobalIteratorAb.html | 597 ++++++++ ...tlass_1_1gemm_1_1IgemmGlobalIteratorAb.png | Bin 0 -> 4485 bytes ...1_1IgemmGlobalLoadTransformer-members.html | 2 +- ...1_1gemm_1_1IgemmGlobalLoadTransformer.html | 2 +- ...ements___01_4_00_01float_01_4-members.html | 2 +- ...00_01kElements___01_4_00_01float_01_4.html | 2 +- ..._1IgemmGlobalStoreTransformer-members.html | 2 +- ..._1gemm_1_1IgemmGlobalStoreTransformer.html | 2 +- ..._t_00_01kElements___01_4_01_4-members.html | 2 +- ..._01int8__t_00_01kElements___01_4_01_4.html | 2 +- ...gemm_1_1IgemmGlobalTileTraits-members.html | 105 ++ ...lass_1_1gemm_1_1IgemmGlobalTileTraits.html | 257 ++++ ...tlass_1_1gemm_1_1IgemmGlobalTileTraits.png | Bin 0 -> 2081 bytes ...balTileTraits_1_1ThreadOffset-members.html | 91 ++ ...IgemmGlobalTileTraits_1_1ThreadOffset.html | 132 ++ ..._1_1IgemmInt8ToFloatConverter-members.html | 2 +- ..._1_1gemm_1_1IgemmInt8ToFloatConverter.html | 2 +- ..._1IgemmSharedStoreTransformer-members.html | 2 +- ..._1gemm_1_1IgemmSharedStoreTransformer.html | 2 +- ...tlass_1_1gemm_1_1IgemmSwizzle-members.html | 2 +- ...structcutlass_1_1gemm_1_1IgemmSwizzle.html | 2 +- ...ass_1_1gemm_1_1IgemmTileTraitsHelperA.html | 14 +- ...lass_1_1gemm_1_1IgemmTileTraitsHelperA.png | Bin 1401 -> 1476 bytes ...emmConfig___00_01Index___01_4-members.html | 99 ++ ...r_00_01GemmConfig___00_01Index___01_4.html | 237 +++ ...or_00_01GemmConfig___00_01Index___01_4.png | Bin 0 -> 1837 bytes ...emmConfig___00_01Index___01_4-members.html | 98 ++ ...r_00_01GemmConfig___00_01Index___01_4.html | 273 ++++ ...ass_1_1gemm_1_1IgemmTileTraitsHelperB.html | 14 +- ...lass_1_1gemm_1_1IgemmTileTraitsHelperB.png | Bin 1400 -> 1476 bytes ...emmConfig___00_01Index___01_4-members.html | 98 ++ ...r_00_01GemmConfig___00_01Index___01_4.html | 273 ++++ ...emmConfig___00_01Index___01_4-members.html | 99 ++ ...r_00_01GemmConfig___00_01Index___01_4.html | 237 +++ ...or_00_01GemmConfig___00_01Index___01_4.png | Bin 0 -> 1727 bytes ...utlass_1_1gemm_1_1IgemmTraits-members.html | 19 +- .../structcutlass_1_1gemm_1_1IgemmTraits.html | 37 +- docs/structcutlass_1_1gemm_1_1IgemmTraits.png | Bin 3096 -> 3065 bytes ..._1_1gemm_1_1IgemmTraitsHelper-members.html | 42 +- ...tcutlass_1_1gemm_1_1IgemmTraitsHelper.html | 256 ++-- ...tcutlass_1_1gemm_1_1IgemmTransformerA.html | 2 +- ...umnMajor_00_01Iterator___01_4-members.html | 2 +- ..._1_1kColumnMajor_00_01Iterator___01_4.html | 2 +- ...RowMajor_00_01Iterator___01_4-members.html | 2 +- ...out_1_1kRowMajor_00_01Iterator___01_4.html | 2 +- ...tcutlass_1_1gemm_1_1IgemmTransformerB.html | 2 +- ...umnMajor_00_01Iterator___01_4-members.html | 2 +- ..._1_1kColumnMajor_00_01Iterator___01_4.html | 2 +- ...RowMajor_00_01Iterator___01_4-members.html | 2 +- ...out_1_1kRowMajor_00_01Iterator___01_4.html | 2 +- ...ructcutlass_1_1gemm_1_1Launch-members.html | 91 ++ docs/structcutlass_1_1gemm_1_1Launch.html | 155 ++ ...unch_3_01Gemm_00_01false_01_4-members.html | 91 ++ ...mm_1_1Launch_3_01Gemm_00_01false_01_4.html | 155 ++ ...lass_1_1gemm_1_1LinearScaling-members.html | 16 +- ...tructcutlass_1_1gemm_1_1LinearScaling.html | 254 ++- ...structcutlass_1_1gemm_1_1LinearScaling.png | Bin 0 -> 1511 bytes ...emm_1_1LinearScalingDevicePtr-members.html | 103 ++ ...ass_1_1gemm_1_1LinearScalingDevicePtr.html | 232 +++ ...lass_1_1gemm_1_1LinearScalingDevicePtr.png | Bin 0 -> 1507 bytes ...mm_1_1LinearScaling_1_1Params-members.html | 6 +- ...ss_1_1gemm_1_1LinearScaling_1_1Params.html | 96 +- ...ructcutlass_1_1gemm_1_1ProjectOperand.html | 2 +- ...rand_1_1kA_00_01Kstrided_01_4-members.html | 2 +- ...1GemmOperand_1_1kA_00_01Kstrided_01_4.html | 2 +- ...rand_1_1kB_00_01Kstrided_01_4-members.html | 2 +- ...1GemmOperand_1_1kB_00_01Kstrided_01_4.html | 2 +- ...mOperand_1_1kC_00_01true_01_4-members.html | 2 +- ..._3_01GemmOperand_1_1kC_00_01true_01_4.html | 2 +- ...mOperand_1_1kD_00_01true_01_4-members.html | 2 +- ..._3_01GemmOperand_1_1kD_00_01true_01_4.html | 2 +- ...ass_1_1gemm_1_1ReshapeThreads-members.html | 2 +- ...ructcutlass_1_1gemm_1_1ReshapeThreads.html | 2 +- ...00_01Threads___00_01true_01_4-members.html | 2 +- ...1Tile___00_01Threads___00_01true_01_4.html | 2 +- ...1gemm_1_1RowMajorBlockSwizzle-members.html | 95 ++ ...tlass_1_1gemm_1_1RowMajorBlockSwizzle.html | 260 ++++ ...utlass_1_1gemm_1_1SgemmConfig-members.html | 59 +- .../structcutlass_1_1gemm_1_1SgemmConfig.html | 163 +- docs/structcutlass_1_1gemm_1_1SgemmConfig.png | Bin 2930 -> 3126 bytes ...lass_1_1gemm_1_1SgemmLBTraits-members.html | 114 ++ ...tructcutlass_1_1gemm_1_1SgemmLBTraits.html | 185 +++ ...structcutlass_1_1gemm_1_1SgemmLBTraits.png | Bin 0 -> 8497 bytes ...utlass_1_1gemm_1_1SgemmTraits-members.html | 19 +- .../structcutlass_1_1gemm_1_1SgemmTraits.html | 37 +- docs/structcutlass_1_1gemm_1_1SgemmTraits.png | Bin 8487 -> 8476 bytes ...s_1_1gemm_1_1SharedLoadStream-members.html | 31 +- ...ctcutlass_1_1gemm_1_1SharedLoadStream.html | 251 ++- ...1_1SharedLoadStream_1_1Params-members.html | 2 +- ...1_1gemm_1_1SharedLoadStream_1_1Params.html | 2 +- ...s_1_1gemm_1_1SharedStreamPair-members.html | 101 ++ ...ctcutlass_1_1gemm_1_1SharedStreamPair.html | 411 +++++ ...1_1SharedStreamPair_1_1Params-members.html | 92 ++ ...1_1gemm_1_1SharedStreamPair_1_1Params.html | 139 ++ ...1SimplifiedGemmEpilogueTraits-members.html | 34 +- ...1gemm_1_1SimplifiedGemmEpilogueTraits.html | 100 +- ..._1gemm_1_1SimplifiedGemmEpilogueTraits.png | Bin 3650 -> 3705 bytes ...1gemm_1_1SimplifiedGemmTraits-members.html | 15 +- ...tlass_1_1gemm_1_1SimplifiedGemmTraits.html | 27 +- ...1_1SimplifiedGemmTraitsHelper-members.html | 6 +- ...1_1gemm_1_1SimplifiedGemmTraitsHelper.html | 26 +- ..._1_1gemm_1_1ThreadMultiplyAdd-members.html | 31 +- ...tcutlass_1_1gemm_1_1ThreadMultiplyAdd.html | 201 +-- ...1half030e27fde4380ad93cd574bc743e0ba3.html | 103 ++ ...1half8ea66703da782e035d986e48031ff835.html | 103 ++ ...__00_01half_00_01half_00_01float_01_4.html | 401 +++++ ...___00_01half_00_01half_00_01half_01_4.html | 402 +++++ ...1int86f255c0f34c1afba22b3a7d64d8f85bf.html | 103 ++ ..._01int8__t_00_01int8__t_00_01int_01_4.html | 401 +++++ ...m_1_1WmmaGemmGlobalIteratorCd-members.html | 82 +- ...s_1_1gemm_1_1WmmaGemmGlobalIteratorCd.html | 717 ++++----- ...ss_1_1gemm_1_1WmmaGemmGlobalIteratorCd.png | Bin 1763 -> 2539 bytes ...mmaGemmGlobalIteratorCdTraits-members.html | 11 +- ...emm_1_1WmmaGemmGlobalIteratorCdTraits.html | 31 +- ...ratorCdTraits_1_1ThreadOffset-members.html | 2 +- ...lobalIteratorCdTraits_1_1ThreadOffset.html | 2 +- ...emmGlobalIteratorCd_1_1Params-members.html | 19 +- ...1_1WmmaGemmGlobalIteratorCd_1_1Params.html | 203 +-- ..._1_1WmmaGemmGlobalIteratorCd_1_1Params.png | Bin 0 -> 1460 bytes ...s_1_1gemm_1_1swizzleDirection-members.html | 93 ++ ...ctcutlass_1_1gemm_1_1swizzleDirection.html | 122 ++ docs/structcutlass_1_1int4__t.html | 92 ++ docs/structcutlass_1_1is__pow2-members.html | 2 +- docs/structcutlass_1_1is__pow2.html | 2 +- docs/structcutlass_1_1log2__down-members.html | 2 +- docs/structcutlass_1_1log2__down.html | 2 +- ..._3_01N_00_011_00_01Count_01_4-members.html | 2 +- ...g2__down_3_01N_00_011_00_01Count_01_4.html | 2 +- docs/structcutlass_1_1log2__up-members.html | 2 +- docs/structcutlass_1_1log2__up.html | 2 +- ..._3_01N_00_011_00_01Count_01_4-members.html | 2 +- ...log2__up_3_01N_00_011_00_01Count_01_4.html | 2 +- ...cutlass_1_1platform_1_1aligned__chunk.html | 2 +- ...1platform_1_1aligned__storage-members.html | 2 +- ...tlass_1_1platform_1_1aligned__storage.html | 2 +- ..._1_1platform_1_1alignment__of-members.html | 2 +- ...tcutlass_1_1platform_1_1alignment__of.html | 2 +- ...tform_1_1alignment__of_1_1pad-members.html | 2 +- ...s_1_1platform_1_1alignment__of_1_1pad.html | 2 +- ..._of_3_01const_01value__t_01_4-members.html | 2 +- ...ignment__of_3_01const_01value__t_01_4.html | 2 +- ...st_01volatile_01value__t_01_4-members.html | 2 +- ..._3_01const_01volatile_01value__t_01_4.html | 2 +- ...lignment__of_3_01double2_01_4-members.html | 2 +- ...orm_1_1alignment__of_3_01double2_01_4.html | 2 +- ...lignment__of_3_01double4_01_4-members.html | 2 +- ...orm_1_1alignment__of_3_01double4_01_4.html | 2 +- ...alignment__of_3_01float4_01_4-members.html | 2 +- ...form_1_1alignment__of_3_01float4_01_4.html | 2 +- ..._1alignment__of_3_01int4_01_4-members.html | 2 +- ...atform_1_1alignment__of_3_01int4_01_4.html | 2 +- ...1alignment__of_3_01long4_01_4-members.html | 2 +- ...tform_1_1alignment__of_3_01long4_01_4.html | 2 +- ...gnment__of_3_01longlong2_01_4-members.html | 2 +- ...m_1_1alignment__of_3_01longlong2_01_4.html | 2 +- ...gnment__of_3_01longlong4_01_4-members.html | 2 +- ...m_1_1alignment__of_3_01longlong4_01_4.html | 2 +- ...1alignment__of_3_01uint4_01_4-members.html | 2 +- ...tform_1_1alignment__of_3_01uint4_01_4.html | 2 +- ...alignment__of_3_01ulong4_01_4-members.html | 2 +- ...form_1_1alignment__of_3_01ulong4_01_4.html | 2 +- ...nment__of_3_01ulonglong2_01_4-members.html | 2 +- ..._1_1alignment__of_3_01ulonglong2_01_4.html | 2 +- ...nment__of_3_01ulonglong4_01_4-members.html | 2 +- ..._1_1alignment__of_3_01ulonglong4_01_4.html | 2 +- ..._3_01volatile_01value__t_01_4-members.html | 2 +- ...ment__of_3_01volatile_01value__t_01_4.html | 2 +- ...1_1platform_1_1bool__constant-members.html | 2 +- ...cutlass_1_1platform_1_1bool__constant.html | 2 +- ...ss_1_1platform_1_1conditional-members.html | 2 +- ...uctcutlass_1_1platform_1_1conditional.html | 2 +- ..._3_01false_00_01T_00_01F_01_4-members.html | 2 +- ...ditional_3_01false_00_01T_00_01F_01_4.html | 2 +- ..._1platform_1_1default__delete-members.html | 2 +- ...utlass_1_1platform_1_1default__delete.html | 2 +- ..._1_1default__delete_3_01T[]_4-members.html | 2 +- ...platform_1_1default__delete_3_01T[]_4.html | 2 +- ...ass_1_1platform_1_1enable__if-members.html | 2 +- ...ructcutlass_1_1platform_1_1enable__if.html | 2 +- ...m_1_1enable__if_3_01false_00_01T_01_4.html | 2 +- ...utlass_1_1platform_1_1greater-members.html | 2 +- .../structcutlass_1_1platform_1_1greater.html | 2 +- ...latform_1_1integral__constant-members.html | 2 +- ...ass_1_1platform_1_1integral__constant.html | 2 +- ...1_1platform_1_1is__arithmetic-members.html | 2 +- ...cutlass_1_1platform_1_1is__arithmetic.html | 2 +- ...s_1_1platform_1_1is__base__of-members.html | 2 +- ...ctcutlass_1_1platform_1_1is__base__of.html | 2 +- ...tform_1_1is__base__of__helper-members.html | 2 +- ...s_1_1platform_1_1is__base__of__helper.html | 2 +- ...is__base__of__helper_1_1dummy-members.html | 2 +- ...form_1_1is__base__of__helper_1_1dummy.html | 2 +- ...atform_1_1is__floating__point-members.html | 2 +- ...ss_1_1platform_1_1is__floating__point.html | 2 +- ..._1platform_1_1is__fundamental-members.html | 2 +- ...utlass_1_1platform_1_1is__fundamental.html | 2 +- ...s_1_1platform_1_1is__integral-members.html | 2 +- ...ctcutlass_1_1platform_1_1is__integral.html | 2 +- ...1_1is__integral_3_01char_01_4-members.html | 2 +- ...latform_1_1is__integral_3_01char_01_4.html | 2 +- ...__integral_3_01const_01T_01_4-members.html | 2 +- ...rm_1_1is__integral_3_01const_01T_01_4.html | 2 +- ...3_01const_01volatile_01T_01_4-members.html | 2 +- ...ntegral_3_01const_01volatile_01T_01_4.html | 2 +- ..._1_1is__integral_3_01int_01_4-members.html | 2 +- ...platform_1_1is__integral_3_01int_01_4.html | 2 +- ...1_1is__integral_3_01long_01_4-members.html | 2 +- ...latform_1_1is__integral_3_01long_01_4.html | 2 +- ...integral_3_01long_01long_01_4-members.html | 2 +- ..._1_1is__integral_3_01long_01long_01_4.html | 2 +- ..._1is__integral_3_01short_01_4-members.html | 2 +- ...atform_1_1is__integral_3_01short_01_4.html | 2 +- ...tegral_3_01signed_01char_01_4-members.html | 2 +- ..._1is__integral_3_01signed_01char_01_4.html | 2 +- ...gral_3_01unsigned_01char_01_4-members.html | 2 +- ...is__integral_3_01unsigned_01char_01_4.html | 2 +- ...egral_3_01unsigned_01int_01_4-members.html | 2 +- ...1is__integral_3_01unsigned_01int_01_4.html | 2 +- ...gral_3_01unsigned_01long_01_4-members.html | 2 +- ...is__integral_3_01unsigned_01long_01_4.html | 2 +- ...01unsigned_01long_01long_01_4-members.html | 2 +- ...egral_3_01unsigned_01long_01long_01_4.html | 2 +- ...ral_3_01unsigned_01short_01_4-members.html | 2 +- ...s__integral_3_01unsigned_01short_01_4.html | 2 +- ...ntegral_3_01volatile_01T_01_4-members.html | 2 +- ...1_1is__integral_3_01volatile_01T_01_4.html | 2 +- ...ss_1_1platform_1_1is__pointer-members.html | 2 +- ...uctcutlass_1_1platform_1_1is__pointer.html | 2 +- ...atform_1_1is__pointer__helper-members.html | 2 +- ...ss_1_1platform_1_1is__pointer__helper.html | 2 +- ...inter__helper_3_01T_01_5_01_4-members.html | 2 +- ..._1is__pointer__helper_3_01T_01_5_01_4.html | 2 +- ...tlass_1_1platform_1_1is__same-members.html | 2 +- ...structcutlass_1_1platform_1_1is__same.html | 2 +- ...1_1is__same_3_01A_00_01A_01_4-members.html | 2 +- ...latform_1_1is__same_3_01A_00_01A_01_4.html | 2 +- ...rm_1_1is__trivially__copyable-members.html | 2 +- ..._1platform_1_1is__trivially__copyable.html | 2 +- ...tlass_1_1platform_1_1is__void-members.html | 2 +- ...structcutlass_1_1platform_1_1is__void.html | 2 +- ...s_1_1platform_1_1is__volatile-members.html | 2 +- ...ctcutlass_1_1platform_1_1is__volatile.html | 2 +- ...olatile_3_01volatile_01T_01_4-members.html | 2 +- ...1_1is__volatile_3_01volatile_01T_01_4.html | 2 +- ...ctcutlass_1_1platform_1_1less-members.html | 2 +- docs/structcutlass_1_1platform_1_1less.html | 2 +- ...ructcutlass_1_1platform_1_1nullptr__t.html | 2 +- ...ctcutlass_1_1platform_1_1plus-members.html | 2 +- docs/structcutlass_1_1platform_1_1plus.html | 2 +- ..._1_1platform_1_1remove__const-members.html | 2 +- ...tcutlass_1_1platform_1_1remove__const.html | 2 +- ...ove__const_3_01const_01T_01_4-members.html | 2 +- ...m_1_1remove__const_3_01const_01T_01_4.html | 2 +- ...ass_1_1platform_1_1remove__cv-members.html | 2 +- ...ructcutlass_1_1platform_1_1remove__cv.html | 2 +- ...1platform_1_1remove__volatile-members.html | 2 +- ...tlass_1_1platform_1_1remove__volatile.html | 2 +- ...olatile_3_01volatile_01T_01_4-members.html | 2 +- ...emove__volatile_3_01volatile_01T_01_4.html | 2 +- docs/structcutlass_1_1sqrt__est-members.html | 2 +- docs/structcutlass_1_1sqrt__est.html | 2 +- docs/structcutlass_1_1uint4__t.html | 92 ++ docs/tensor__ref_8h.html | 19 +- docs/tensor__ref_8h_source.html | 102 +- docs/tensor__ref__collection_8h.html | 115 ++ docs/tensor__ref__collection_8h_source.html | 146 ++ docs/tensor__view_8h.html | 15 +- docs/tensor__view_8h_source.html | 82 +- docs/thread__multiply__add_8h.html | 6 +- docs/thread__multiply__add_8h_source.html | 33 +- docs/threadblock__swizzle_8h.html | 126 ++ docs/threadblock__swizzle_8h_source.html | 120 ++ docs/tile__allocation_8h.html | 114 ++ docs/tile__allocation_8h_source.html | 121 ++ docs/tile__coord_8h.html | 107 ++ docs/tile__coord_8h_source.html | 128 ++ docs/tile__iterator_8h.html | 35 +- docs/tile__iterator_8h_source.html | 353 +++-- docs/tile__stream_8h.html | 130 ++ docs/tile__stream_8h_source.html | 162 ++ docs/tile__traits__standard_8h.html | 4 +- docs/tile__traits__standard_8h_source.html | 13 +- docs/unioncutlass_1_1Vector-members.html | 6 +- docs/unioncutlass_1_1Vector.html | 26 +- ..._01bin1__t_00_01kLanes___01_4-members.html | 98 ++ ...Vector_3_01bin1__t_00_01kLanes___01_4.html | 329 ++++ ..._1Vector_3_01half_00_011_01_4-members.html | 97 ++ ...utlass_1_1Vector_3_01half_00_011_01_4.html | 279 ++++ ...r_3_01half_00_01kLanes___01_4-members.html | 10 +- ...1_1Vector_3_01half_00_01kLanes___01_4.html | 68 +- ..._01int4__t_00_01kLanes___01_4-members.html | 98 ++ ...Vector_3_01int4__t_00_01kLanes___01_4.html | 329 ++++ ...01uint4__t_00_01kLanes___01_4-members.html | 98 ++ ...ector_3_01uint4__t_00_01kLanes___01_4.html | 329 ++++ ...Traits_1_1StreamSharedStorage-members.html | 10 +- ...EpilogueTraits_1_1StreamSharedStorage.html | 30 +- ..._1GemmTraits_1_1SharedStorage-members.html | 2 +- ..._1gemm_1_1GemmTraits_1_1SharedStorage.html | 6 +- docs/vector_8h.html | 34 +- docs/vector_8h_source.html | 90 +- docs/wmma__gemm__epilogue__traits_8h.html | 4 +- ...mma__gemm__epilogue__traits_8h_source.html | 14 +- docs/wmma__gemm__global__tile_8h.html | 4 +- docs/wmma__gemm__global__tile_8h_source.html | 92 +- docs/wmma__gemm__multiply__add_8h.html | 4 +- docs/wmma__gemm__multiply__add_8h_source.html | 11 +- docs/wmma__gemm__shared__tile_8h.html | 4 +- docs/wmma__gemm__shared__tile_8h_source.html | 20 +- docs/wmma__gemm__traits_8h.html | 4 +- docs/wmma__gemm__traits_8h_source.html | 24 +- docs/wmma__matrix_8h.html | 2 +- docs/wmma__matrix_8h_source.html | 12 +- docs/zip__fragment_8h.html | 127 ++ docs/zip__fragment_8h_source.html | 114 ++ docs/zip__tensor__ref_8h.html | 116 ++ docs/zip__tensor__ref_8h_source.html | 99 ++ docs/zip__tile__iterator_8h.html | 113 ++ docs/zip__tile__iterator_8h_source.html | 139 ++ 1121 files changed, 64211 insertions(+), 12918 deletions(-) create mode 100644 docs/classcutlass_1_1TensorRef.png create mode 100644 docs/classcutlass_1_1TensorRefArray_1_1ConstIterator-members.html create mode 100644 docs/classcutlass_1_1TensorRefArray_1_1ConstIterator.html create mode 100644 docs/classcutlass_1_1TensorRefBatchStrided_1_1ConstIterator-members.html create mode 100644 docs/classcutlass_1_1TensorRefBatchStrided_1_1ConstIterator.html create mode 100644 docs/classcutlass_1_1TensorRef_3_01Storage___00_01Rank___00_01MapFunc___00_011_00_01Index___00_01LongIndex___01_4-members.html create mode 100644 docs/classcutlass_1_1TensorRef_3_01Storage___00_01Rank___00_01MapFunc___00_011_00_01Index___00_01LongIndex___01_4.html create mode 100644 docs/classcutlass_1_1ZipTileIterator-members.html create mode 100644 docs/classcutlass_1_1ZipTileIterator.html create mode 100644 docs/classcutlass_1_1detail_1_1ScalarOrPointer-members.html create mode 100644 docs/classcutlass_1_1detail_1_1ScalarOrPointer.html create mode 100644 docs/classcutlass_1_1gemm_1_1LinearScalingDevicePtr_1_1Params-members.html create mode 100644 docs/classcutlass_1_1gemm_1_1LinearScalingDevicePtr_1_1Params.html create mode 100644 docs/classcutlass_1_1platform_1_1complex-members.html create mode 100644 docs/classcutlass_1_1platform_1_1complex.html create mode 100644 docs/complex_8h.html create mode 100644 docs/complex_8h_source.html create mode 100644 docs/fp16__sgemm__multiply__add_8h.html create mode 100644 docs/fp16__sgemm__multiply__add_8h_source.html create mode 100644 docs/fp16__sgemm__traits_8h.html create mode 100644 docs/fp16__sgemm__traits_8h_source.html create mode 100644 docs/functions_func_k.html create mode 100644 docs/functions_func_n.html create mode 100644 docs/functions_func_z.html create mode 100644 docs/functions_type_k.html create mode 100644 docs/functions_vars_o.html create mode 100644 docs/functions_z.html create mode 100644 docs/gemm__config_8h.html create mode 100644 docs/gemm__config_8h_source.html create mode 100644 docs/gemm__coord_8h.html create mode 100644 docs/gemm__coord_8h_source.html create mode 100644 docs/gemm__desc_8h.html create mode 100644 docs/gemm__desc_8h_source.html create mode 100644 docs/gemm__stream__pair_8h.html create mode 100644 docs/gemm__stream__pair_8h_source.html create mode 100644 docs/group__IdentityBlockSwizzle.html create mode 100644 docs/kernel__launch_8h.html create mode 100644 docs/kernel__launch_8h_source.html create mode 100644 docs/linear__scaling__device__ptr_8h.html create mode 100644 docs/linear__scaling__device__ptr_8h_source.html create mode 100644 docs/namespacecutlass_1_1detail.html create mode 100644 docs/numeric__types_8h.html create mode 100644 docs/numeric__types_8h_source.html create mode 100644 docs/scalar__or__pointer_8h.html create mode 100644 docs/scalar__or__pointer_8h_source.html create mode 100644 docs/search/all_18.html create mode 100644 docs/search/all_18.js create mode 100644 docs/search/classes_13.html create mode 100644 docs/search/classes_13.js create mode 100644 docs/search/classes_14.html create mode 100644 docs/search/classes_14.js create mode 100644 docs/search/enumvalues_4.html create mode 100644 docs/search/enumvalues_4.js create mode 100644 docs/search/enumvalues_5.html create mode 100644 docs/search/enumvalues_5.js create mode 100644 docs/search/files_10.html create mode 100644 docs/search/files_10.js create mode 100644 docs/search/files_f.html create mode 100644 docs/search/files_f.js create mode 100644 docs/search/functions_15.html create mode 100644 docs/search/functions_15.js create mode 100644 docs/search/functions_16.html create mode 100644 docs/search/functions_16.js create mode 100644 docs/search/functions_17.html create mode 100644 docs/search/functions_17.js create mode 100644 docs/search/groups_4.html create mode 100644 docs/search/groups_4.js create mode 100644 docs/search/typedefs_12.html create mode 100644 docs/search/typedefs_12.js create mode 100644 docs/search/variables_11.html create mode 100644 docs/search/variables_11.js create mode 100644 docs/structDebugType.html create mode 100644 docs/structDebugValue.html create mode 100644 docs/structcutlass_1_1DumpType.html create mode 100644 docs/structcutlass_1_1FragmentElementType-members.html create mode 100644 docs/structcutlass_1_1FragmentElementType.html create mode 100644 docs/structcutlass_1_1IdentityTensorMapFunc-members.html create mode 100644 docs/structcutlass_1_1IdentityTensorMapFunc.html create mode 100644 docs/structcutlass_1_1KernelLaunchConfiguration-members.html create mode 100644 docs/structcutlass_1_1KernelLaunchConfiguration.html create mode 100644 docs/structcutlass_1_1KernelLaunchConfiguration.png create mode 100644 docs/structcutlass_1_1Load_3_01Scalar___00_01kAccessSize_00_01Memory___00_01FragmentElementType_1_1kS1b28106546bd22002cd52e60197548ce.html create mode 100644 docs/structcutlass_1_1Load_3_01Scalar___00_01kAccessSize_00_01Memory___00_01FragmentElementType_1_1kS3de71ba9ef724a37fb1cf315da8bf9b5.html create mode 100644 docs/structcutlass_1_1Load_3_01Scalar___00_01kAccessSize_00_01Memory___00_01FragmentElementType_1_1kS840de374cd1e3ee3a4aa6a70a9fa83d1.html create mode 100644 docs/structcutlass_1_1Load_3_01Scalar___00_01kAccessSize_00_01Memory___00_01FragmentElementType_1_1kS91ff65dd36f118bca4542df6128d6d99.html create mode 100644 docs/structcutlass_1_1Load_3_01Scalar___00_01kAccessSize_00_01Memory___00_01FragmentElementType_1_1kS966cdf9ba449f056a92458bf87878053.html create mode 100644 docs/structcutlass_1_1Load_3_01Scalar___00_01kAccessSize_00_01Memory___00_01FragmentElementType_1_1kSc8e65d97c8f615e98ac2e03f0bcd236e.html create mode 100644 docs/structcutlass_1_1Load_3_01Scalar___00_01kAccessSize_00_01Memory___00_01FragmentElementType_1_1kSf1f63874ddc1302ed9e60e0478e1a8ad.html create mode 100644 docs/structcutlass_1_1Load_3_01Scalar___00_01kAccessSize_00_01Memory___00_01FragmentElementType_1_1kSf2952bcb62d20c76c595eac4c59f7239.html create mode 100644 docs/structcutlass_1_1Load_3_01Scalar___00_01kAccessSize_00_01Memory___00_01FragmentElementType_1_1kW847237836867f08e1121b00b7e44d8ae.html create mode 100644 docs/structcutlass_1_1Load_3_01Scalar___00_01kAccessSize_00_01Memory___00_01FragmentElementType_1_1kWe58d0048e24352beeec002fd483c53b2.html create mode 100644 docs/structcutlass_1_1Load_3_01Vector_3_01bin1__t_00_0132_01_4_00_01kAccessSize_00_01Memory___00_01Fr16ad55dd7bca84e6a7fc608c26eee889.html create mode 100644 docs/structcutlass_1_1Load_3_01Vector_3_01bin1__t_00_0132_01_4_00_01kAccessSize_00_01Memory___00_01Fra240d9c28383cd3945277ec3a927c538.html create mode 100644 docs/structcutlass_1_1Load_3_01Vector_3_01int4__t_00_018_01_4_00_01kAccessSize_00_01Memory___00_01Fra03475418f27732d726d18feb23feeed2.html create mode 100644 docs/structcutlass_1_1Load_3_01Vector_3_01int4__t_00_018_01_4_00_01kAccessSize_00_01Memory___00_01Fra934529165fa8fecdd392b5302d25ef26.html create mode 100644 docs/structcutlass_1_1Load_3_01Vector_3_01uint4__t_00_018_01_4_00_01kAccessSize_00_01Memory___00_01Fr23088b868c6eeec9377c46892553686f.html create mode 100644 docs/structcutlass_1_1Load_3_01Vector_3_01uint4__t_00_018_01_4_00_01kAccessSize_00_01Memory___00_01Frc8b0aa6ddd9b4317158c26574a6881de.html create mode 100644 docs/structcutlass_1_1Load_3_01double_00_012_00_01Memory___00_01FragmentElementType_1_1kScalar_00_01d9e675253ca19588f1ae4bd898579523f.html create mode 100644 docs/structcutlass_1_1Load_3_01double_00_012_00_01Memory___00_01FragmentElementType_1_1kScalar_00_01double_00_01kStride_00_0116_01_4.html create mode 100644 docs/structcutlass_1_1MatrixCoord-members.html create mode 100644 docs/structcutlass_1_1MatrixCoord.html create mode 100644 docs/structcutlass_1_1MatrixCoord.png create mode 100644 docs/structcutlass_1_1MatrixLayout_1_1ColumnMajor-members.html create mode 100644 docs/structcutlass_1_1MatrixLayout_1_1ColumnMajor.html create mode 100644 docs/structcutlass_1_1MatrixLayout_1_1ColumnMajorBlockLinear-members.html create mode 100644 docs/structcutlass_1_1MatrixLayout_1_1ColumnMajorBlockLinear.html create mode 100644 docs/structcutlass_1_1MatrixLayout_1_1ColumnMajorInterleaved-members.html create mode 100644 docs/structcutlass_1_1MatrixLayout_1_1ColumnMajorInterleaved.html create mode 100644 docs/structcutlass_1_1MatrixLayout_1_1ContiguousLayout-members.html create mode 100644 docs/structcutlass_1_1MatrixLayout_1_1ContiguousLayout.html create mode 100644 docs/structcutlass_1_1MatrixLayout_1_1RowMajor-members.html create mode 100644 docs/structcutlass_1_1MatrixLayout_1_1RowMajor.html create mode 100644 docs/structcutlass_1_1MatrixLayout_1_1RowMajorBlockLinear-members.html create mode 100644 docs/structcutlass_1_1MatrixLayout_1_1RowMajorBlockLinear.html create mode 100644 docs/structcutlass_1_1MatrixLayout_1_1RowMajorInterleaved-members.html create mode 100644 docs/structcutlass_1_1MatrixLayout_1_1RowMajorInterleaved.html create mode 100644 docs/structcutlass_1_1MatrixTransform-members.html create mode 100644 docs/structcutlass_1_1MatrixTransform.html create mode 100644 docs/structcutlass_1_1Max-members.html create mode 100644 docs/structcutlass_1_1Max.html create mode 100644 docs/structcutlass_1_1Min-members.html create mode 100644 docs/structcutlass_1_1Min.html create mode 100644 docs/structcutlass_1_1PredicatedTileLoadStream-members.html create mode 100644 docs/structcutlass_1_1PredicatedTileLoadStream.html create mode 100644 docs/structcutlass_1_1PredicatedTileLoadStream.png create mode 100644 docs/structcutlass_1_1PredicatedTileStoreStream-members.html create mode 100644 docs/structcutlass_1_1PredicatedTileStoreStream.html create mode 100644 docs/structcutlass_1_1PredicatedTileStoreStream.png create mode 100644 docs/structcutlass_1_1RegularTilePredicateFunctor-members.html create mode 100644 docs/structcutlass_1_1RegularTilePredicateFunctor.html create mode 100644 docs/structcutlass_1_1ScalarIO-members.html create mode 100644 docs/structcutlass_1_1ScalarIO.html create mode 100644 docs/structcutlass_1_1ShapeDivCeiling-members.html create mode 100644 docs/structcutlass_1_1ShapeDivCeiling.html create mode 100644 docs/structcutlass_1_1Store_3_01Scalar___00_01kAccessSize_00_01Memory___00_01FragmentElementType_1_1k004b304998a534d76357f834068909f8.html create mode 100644 docs/structcutlass_1_1Store_3_01Scalar___00_01kAccessSize_00_01Memory___00_01FragmentElementType_1_1k12f5c8a016a307e76de374322fc00a66.html create mode 100644 docs/structcutlass_1_1Store_3_01Scalar___00_01kAccessSize_00_01Memory___00_01FragmentElementType_1_1k220d5790f803f10840e2a92fb9a51dac.html create mode 100644 docs/structcutlass_1_1Store_3_01Scalar___00_01kAccessSize_00_01Memory___00_01FragmentElementType_1_1k28cc0b88a16efca73d258128312d2a7e.html create mode 100644 docs/structcutlass_1_1Store_3_01Scalar___00_01kAccessSize_00_01Memory___00_01FragmentElementType_1_1k40d038d4bce377843c21a56ebf97d011.html create mode 100644 docs/structcutlass_1_1Store_3_01Scalar___00_01kAccessSize_00_01Memory___00_01FragmentElementType_1_1k60eedca420c41e94fd40b41299967ef2.html create mode 100644 docs/structcutlass_1_1Store_3_01Scalar___00_01kAccessSize_00_01Memory___00_01FragmentElementType_1_1k775a1d27affec5236489735ed4503c92.html create mode 100644 docs/structcutlass_1_1Store_3_01Scalar___00_01kAccessSize_00_01Memory___00_01FragmentElementType_1_1k84da7dcd68ee74b8d2bdb67885b0ca56.html create mode 100644 docs/structcutlass_1_1Store_3_01Scalar___00_01kAccessSize_00_01Memory___00_01FragmentElementType_1_1kd9a7e85f80a21c504388612a60462417.html create mode 100644 docs/structcutlass_1_1Store_3_01Scalar___00_01kAccessSize_00_01Memory___00_01FragmentElementType_1_1ke6d73d34fa7b5254cf828804a19842e1.html create mode 100644 docs/structcutlass_1_1Store_3_01double_00_012_00_01Memory___00_01FragmentElementType_1_1kScalar_00_013d38935f41bf709e067932b9e042255a.html create mode 100644 docs/structcutlass_1_1Store_3_01double_00_012_00_01Memory___00_01FragmentElementType_1_1kScalar_00_0160391c6be5cb1d3f99e012a6a18e486d.html create mode 100644 docs/structcutlass_1_1TensorRefArray-members.html create mode 100644 docs/structcutlass_1_1TensorRefArray.html create mode 100644 docs/structcutlass_1_1TensorRefBatchStrided-members.html create mode 100644 docs/structcutlass_1_1TensorRefBatchStrided.html create mode 100644 docs/structcutlass_1_1TensorRefBatchStrided.png create mode 100644 docs/structcutlass_1_1TensorRef_3_01Storage___00_01Rank___00_01MapFunc___00_011_00_01Index___00_01LongIndex___01_4_1_1StrideVector.html create mode 100644 docs/structcutlass_1_1TileAllocation-members.html create mode 100644 docs/structcutlass_1_1TileAllocation.html create mode 100644 docs/structcutlass_1_1TileCoord-members.html create mode 100644 docs/structcutlass_1_1TileCoord.html create mode 100644 docs/structcutlass_1_1TileCoord.png create mode 100644 docs/structcutlass_1_1TileLoadStream-members.html create mode 100644 docs/structcutlass_1_1TileLoadStream.html create mode 100644 docs/structcutlass_1_1TileLoadStream.png create mode 100644 docs/structcutlass_1_1TileLoadStream_1_1Params-members.html create mode 100644 docs/structcutlass_1_1TileLoadStream_1_1Params.html create mode 100644 docs/structcutlass_1_1TileLoadStream_1_1PredicateVector.html create mode 100644 docs/structcutlass_1_1TileStoreStream-members.html create mode 100644 docs/structcutlass_1_1TileStoreStream.html create mode 100644 docs/structcutlass_1_1TileStoreStream.png create mode 100644 docs/structcutlass_1_1TileStoreStream_1_1Params-members.html create mode 100644 docs/structcutlass_1_1TileStoreStream_1_1Params.html create mode 100644 docs/structcutlass_1_1TileStoreStream_1_1PredicateVector.html create mode 100644 docs/structcutlass_1_1Vectorize_3_01Vector_3_01bin1__t_00_0132_01_4_00_01kLanes___01_4-members.html create mode 100644 docs/structcutlass_1_1Vectorize_3_01Vector_3_01bin1__t_00_0132_01_4_00_01kLanes___01_4.html create mode 100644 docs/structcutlass_1_1Vectorize_3_01Vector_3_01int4__t_00_018_01_4_00_01kLanes___01_4-members.html create mode 100644 docs/structcutlass_1_1Vectorize_3_01Vector_3_01int4__t_00_018_01_4_00_01kLanes___01_4.html create mode 100644 docs/structcutlass_1_1Vectorize_3_01Vector_3_01uint4__t_00_018_01_4_00_01kLanes___01_4-members.html create mode 100644 docs/structcutlass_1_1Vectorize_3_01Vector_3_01uint4__t_00_018_01_4_00_01kLanes___01_4.html create mode 100644 docs/structcutlass_1_1ZipConvert-members.html create mode 100644 docs/structcutlass_1_1ZipConvert.html create mode 100644 docs/structcutlass_1_1ZipFragment-members.html create mode 100644 docs/structcutlass_1_1ZipFragment.html create mode 100644 docs/structcutlass_1_1ZipTensorRef-members.html create mode 100644 docs/structcutlass_1_1ZipTensorRef.html create mode 100644 docs/structcutlass_1_1ZipTileAllocation-members.html create mode 100644 docs/structcutlass_1_1ZipTileAllocation.html create mode 100644 docs/structcutlass_1_1ZipTileIterator_1_1Params-members.html create mode 100644 docs/structcutlass_1_1ZipTileIterator_1_1Params.html create mode 100644 docs/structcutlass_1_1bin1__t.html create mode 100644 docs/structcutlass_1_1gemm_1_1ColumnMajorBlockSwizzle-members.html create mode 100644 docs/structcutlass_1_1gemm_1_1ColumnMajorBlockSwizzle.html create mode 100644 docs/structcutlass_1_1gemm_1_1Fp16SgemmConfig-members.html create mode 100644 docs/structcutlass_1_1gemm_1_1Fp16SgemmConfig.html create mode 100644 docs/structcutlass_1_1gemm_1_1Fp16SgemmConfig.png create mode 100644 docs/structcutlass_1_1gemm_1_1Fp16SgemmSgemmTraits-members.html create mode 100644 docs/structcutlass_1_1gemm_1_1Fp16SgemmSgemmTraits.html create mode 100644 docs/structcutlass_1_1gemm_1_1Fp16SgemmSgemmTraits.png create mode 100644 docs/structcutlass_1_1gemm_1_1FragmentMultiplyAdd_3_01half_00_01half_00_01true_01_4-members.html create mode 100644 docs/structcutlass_1_1gemm_1_1FragmentMultiplyAdd_3_01half_00_01half_00_01true_01_4.html create mode 100644 docs/structcutlass_1_1gemm_1_1GemmCoord-members.html create mode 100644 docs/structcutlass_1_1gemm_1_1GemmCoord.html create mode 100644 docs/structcutlass_1_1gemm_1_1GemmCoord.png create mode 100644 docs/structcutlass_1_1gemm_1_1GemmGlobalIteratorCd_1_1Params.png create mode 100644 docs/structcutlass_1_1gemm_1_1GemmTraits_1_1Params.png create mode 100644 docs/structcutlass_1_1gemm_1_1GlobalLoadStreamPair-members.html create mode 100644 docs/structcutlass_1_1gemm_1_1GlobalLoadStreamPair.html create mode 100644 docs/structcutlass_1_1gemm_1_1GlobalLoadStreamPair_1_1Params-members.html create mode 100644 docs/structcutlass_1_1gemm_1_1GlobalLoadStreamPair_1_1Params.html create mode 100644 docs/structcutlass_1_1gemm_1_1GlobalLoadStreamPair_1_1SharedStorage-members.html create mode 100644 docs/structcutlass_1_1gemm_1_1GlobalLoadStreamPair_1_1SharedStorage.html create mode 100644 docs/structcutlass_1_1gemm_1_1GlobalLoadStream_1_1Params-members.html create mode 100644 docs/structcutlass_1_1gemm_1_1GlobalLoadStream_1_1Params.html create mode 100644 docs/structcutlass_1_1gemm_1_1GlobalLoadStream_1_1SharedStorage.html create mode 100644 docs/structcutlass_1_1gemm_1_1IgemmConfig_3_01OutputTile___00_01int8__t_00_01ThreadGemmShape___01_4-members.html create mode 100644 docs/structcutlass_1_1gemm_1_1IgemmConfig_3_01OutputTile___00_01int8__t_00_01ThreadGemmShape___01_4.html create mode 100644 docs/structcutlass_1_1gemm_1_1IgemmConfig_3_01OutputTile___00_01int8__t_00_01ThreadGemmShape___01_4.png create mode 100644 docs/structcutlass_1_1gemm_1_1IgemmGlobalIteratorAb-members.html create mode 100644 docs/structcutlass_1_1gemm_1_1IgemmGlobalIteratorAb.html create mode 100644 docs/structcutlass_1_1gemm_1_1IgemmGlobalIteratorAb.png create mode 100644 docs/structcutlass_1_1gemm_1_1IgemmGlobalTileTraits-members.html create mode 100644 docs/structcutlass_1_1gemm_1_1IgemmGlobalTileTraits.html create mode 100644 docs/structcutlass_1_1gemm_1_1IgemmGlobalTileTraits.png create mode 100644 docs/structcutlass_1_1gemm_1_1IgemmGlobalTileTraits_1_1ThreadOffset-members.html create mode 100644 docs/structcutlass_1_1gemm_1_1IgemmGlobalTileTraits_1_1ThreadOffset.html create mode 100644 docs/structcutlass_1_1gemm_1_1IgemmTileTraitsHelperA_3_01MatrixLayout_1_1kColumnMajor_00_01GemmConfig___00_01Index___01_4-members.html create mode 100644 docs/structcutlass_1_1gemm_1_1IgemmTileTraitsHelperA_3_01MatrixLayout_1_1kColumnMajor_00_01GemmConfig___00_01Index___01_4.html create mode 100644 docs/structcutlass_1_1gemm_1_1IgemmTileTraitsHelperA_3_01MatrixLayout_1_1kColumnMajor_00_01GemmConfig___00_01Index___01_4.png create mode 100644 docs/structcutlass_1_1gemm_1_1IgemmTileTraitsHelperA_3_01MatrixLayout_1_1kRowMajor_00_01GemmConfig___00_01Index___01_4-members.html create mode 100644 docs/structcutlass_1_1gemm_1_1IgemmTileTraitsHelperA_3_01MatrixLayout_1_1kRowMajor_00_01GemmConfig___00_01Index___01_4.html create mode 100644 docs/structcutlass_1_1gemm_1_1IgemmTileTraitsHelperB_3_01MatrixLayout_1_1kColumnMajor_00_01GemmConfig___00_01Index___01_4-members.html create mode 100644 docs/structcutlass_1_1gemm_1_1IgemmTileTraitsHelperB_3_01MatrixLayout_1_1kColumnMajor_00_01GemmConfig___00_01Index___01_4.html create mode 100644 docs/structcutlass_1_1gemm_1_1IgemmTileTraitsHelperB_3_01MatrixLayout_1_1kRowMajor_00_01GemmConfig___00_01Index___01_4-members.html create mode 100644 docs/structcutlass_1_1gemm_1_1IgemmTileTraitsHelperB_3_01MatrixLayout_1_1kRowMajor_00_01GemmConfig___00_01Index___01_4.html create mode 100644 docs/structcutlass_1_1gemm_1_1IgemmTileTraitsHelperB_3_01MatrixLayout_1_1kRowMajor_00_01GemmConfig___00_01Index___01_4.png create mode 100644 docs/structcutlass_1_1gemm_1_1Launch-members.html create mode 100644 docs/structcutlass_1_1gemm_1_1Launch.html create mode 100644 docs/structcutlass_1_1gemm_1_1Launch_3_01Gemm_00_01false_01_4-members.html create mode 100644 docs/structcutlass_1_1gemm_1_1Launch_3_01Gemm_00_01false_01_4.html create mode 100644 docs/structcutlass_1_1gemm_1_1LinearScaling.png create mode 100644 docs/structcutlass_1_1gemm_1_1LinearScalingDevicePtr-members.html create mode 100644 docs/structcutlass_1_1gemm_1_1LinearScalingDevicePtr.html create mode 100644 docs/structcutlass_1_1gemm_1_1LinearScalingDevicePtr.png create mode 100644 docs/structcutlass_1_1gemm_1_1RowMajorBlockSwizzle-members.html create mode 100644 docs/structcutlass_1_1gemm_1_1RowMajorBlockSwizzle.html create mode 100644 docs/structcutlass_1_1gemm_1_1SgemmLBTraits-members.html create mode 100644 docs/structcutlass_1_1gemm_1_1SgemmLBTraits.html create mode 100644 docs/structcutlass_1_1gemm_1_1SgemmLBTraits.png create mode 100644 docs/structcutlass_1_1gemm_1_1SharedStreamPair-members.html create mode 100644 docs/structcutlass_1_1gemm_1_1SharedStreamPair.html create mode 100644 docs/structcutlass_1_1gemm_1_1SharedStreamPair_1_1Params-members.html create mode 100644 docs/structcutlass_1_1gemm_1_1SharedStreamPair_1_1Params.html create mode 100644 docs/structcutlass_1_1gemm_1_1ThreadMultiplyAdd_3_01ThreadGemmShape___00_01ThreadsPerWarp___00_01half030e27fde4380ad93cd574bc743e0ba3.html create mode 100644 docs/structcutlass_1_1gemm_1_1ThreadMultiplyAdd_3_01ThreadGemmShape___00_01ThreadsPerWarp___00_01half8ea66703da782e035d986e48031ff835.html create mode 100644 docs/structcutlass_1_1gemm_1_1ThreadMultiplyAdd_3_01ThreadGemmShape___00_01ThreadsPerWarp___00_01half_00_01half_00_01float_01_4.html create mode 100644 docs/structcutlass_1_1gemm_1_1ThreadMultiplyAdd_3_01ThreadGemmShape___00_01ThreadsPerWarp___00_01half_00_01half_00_01half_01_4.html create mode 100644 docs/structcutlass_1_1gemm_1_1ThreadMultiplyAdd_3_01ThreadGemmShape___00_01ThreadsPerWarp___00_01int86f255c0f34c1afba22b3a7d64d8f85bf.html create mode 100644 docs/structcutlass_1_1gemm_1_1ThreadMultiplyAdd_3_01ThreadGemmShape___00_01ThreadsPerWarp___00_01int8__t_00_01int8__t_00_01int_01_4.html create mode 100644 docs/structcutlass_1_1gemm_1_1WmmaGemmGlobalIteratorCd_1_1Params.png create mode 100644 docs/structcutlass_1_1gemm_1_1swizzleDirection-members.html create mode 100644 docs/structcutlass_1_1gemm_1_1swizzleDirection.html create mode 100644 docs/structcutlass_1_1int4__t.html create mode 100644 docs/structcutlass_1_1uint4__t.html create mode 100644 docs/tensor__ref__collection_8h.html create mode 100644 docs/tensor__ref__collection_8h_source.html create mode 100644 docs/threadblock__swizzle_8h.html create mode 100644 docs/threadblock__swizzle_8h_source.html create mode 100644 docs/tile__allocation_8h.html create mode 100644 docs/tile__allocation_8h_source.html create mode 100644 docs/tile__coord_8h.html create mode 100644 docs/tile__coord_8h_source.html create mode 100644 docs/tile__stream_8h.html create mode 100644 docs/tile__stream_8h_source.html create mode 100644 docs/unioncutlass_1_1Vector_3_01bin1__t_00_01kLanes___01_4-members.html create mode 100644 docs/unioncutlass_1_1Vector_3_01bin1__t_00_01kLanes___01_4.html create mode 100644 docs/unioncutlass_1_1Vector_3_01half_00_011_01_4-members.html create mode 100644 docs/unioncutlass_1_1Vector_3_01half_00_011_01_4.html create mode 100644 docs/unioncutlass_1_1Vector_3_01int4__t_00_01kLanes___01_4-members.html create mode 100644 docs/unioncutlass_1_1Vector_3_01int4__t_00_01kLanes___01_4.html create mode 100644 docs/unioncutlass_1_1Vector_3_01uint4__t_00_01kLanes___01_4-members.html create mode 100644 docs/unioncutlass_1_1Vector_3_01uint4__t_00_01kLanes___01_4.html create mode 100644 docs/zip__fragment_8h.html create mode 100644 docs/zip__fragment_8h_source.html create mode 100644 docs/zip__tensor__ref_8h.html create mode 100644 docs/zip__tensor__ref_8h_source.html create mode 100644 docs/zip__tile__iterator_8h.html create mode 100644 docs/zip__tile__iterator_8h_source.html diff --git a/docs/annotated.html b/docs/annotated.html index e6c405d5..da54a8ee 100644 --- a/docs/annotated.html +++ b/docs/annotated.html @@ -74,303 +74,368 @@ $(function() {
Here are the classes, structs, unions and interfaces with brief descriptions:
[detail level 1234]
- - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + +
 ▼Ncutlass
 ▶Ngemm
 ▶Nplatform
 CAlignedStruct
 CComputeOffsetFromShapeCompute the offset for the given coordinates in a cube
 CComputeOffsetFromShape< Shape< 1, kSh_, kSw_, 1 > >Compute the offset for the given coordinates in a cube with one channel and a depth of 1
 CComputeOffsetFromShape< Shape< 1, kSh_, kSw_, kSc_ > >Compute the offset for the given coordinates in a cube with a depth of 1
 CComputeOffsetFromStridesCompute the offset for the given coordinates in a cube
 CComputeOffsetFromStrides< Shape< 1, S_h_, S_w_, 1 > >Compute the offset for the given coordinates in a cube with one channel and a depth of 1
 CComputeOffsetFromStrides< Shape< 1, S_h_, S_w_, S_c_ > >Compute the offset for the given coordinates in a cube with a depth of 1
 CComputeThreadOffsetFromStridesDecompose threadId.x into coordinate of a cube whose dimensions are specified by Threads_. Afterwards compute the offset of those coordinates using Strides_
 CComputeThreadOffsetFromStrides< Shape< 1, T_h_, T_w_, 1 >, Shape< 1, S_h_, S_w_, 1 > >Specialization for D=1 and C=1
 CComputeThreadOffsetFromStrides< Shape< 1, T_h_, T_w_, T_c_ >, Shape< 1, S_h_, S_w_, S_c_ > >Specialization for D=1
 CConstPredicateTileAdapterAdapter to enable random access to predicates via logical coordinate within a tile
 CConvert
 CConvert< Fragment< InputScalar_, kScalars_ >, Fragment< OutputScalar_, kScalars_ > >
 CCoordStatically-sized array specifying Coords within a tensor
 CCopy
 Cdivide_assert
 CExtentReturns the extent of a scalar or vector
 CExtent< Vector< T, Lanes > >Returns the number of lanes of a vector if need be
 CExtent< Vector< T, Lanes > const >Returns the number of lanes of a vector if need be
 CFragmentA template defining Fragment Concept
 CFragmentConstIterator
 CFragmentIteratorA template defining Fragment Iterator Concept
 CFragmentLoad
 CFragmentLoad< IteratorFragment::kScalar, kAccessSize, Scalar_, Memory_, FragmentElement_, kStride >
 CFragmentLoad< IteratorFragment::kWmmaMatrix, kAccessSize, Scalar_, Memory_, FragmentElement_, kStride >
 CFragmentStore
 CFragmentStore< IteratorFragment::kScalar, kAccessSize, Scalar_, Memory_, FragmentElement_, kStride >
 CFragmentStore< IteratorFragment::kWmmaMatrix, kAccessSize, Scalar_, Memory_, FragmentElement_, kStride >
 CGemmOperandGemm operand - D = A * B + C
 CIdentityDescribes identity elements
 Cis_pow2
 CIteratorAdvanceSpecifies dimension in which post-increment accesses advance
 CIteratorFragmentSpecifies whether iterator storage fragment consists of Scalar values or WMMA matrix
 CLoad
 CLoad< double, 2, Memory_, true, 16 >
 CLoad< Scalar_, Lanes_, Memory_, true, 16 >
 CLoad< Scalar_, Lanes_, Memory_, true, 4 >
 CLoad< Scalar_, Lanes_, Memory_, true, 8 >
 Clog2_down
 Clog2_down< N, 1, Count >
 Clog2_up
 Clog2_up< N, 1, Count >
 CMatrixLayoutDescribes layouts of matrices
 CMemorySpaceEnum to specify which memory space data resides in
 CPredicateTileAdapterAdapter to enable random access to predicates via logical coordinate within a tile
 ▶CPredicateVectorStatically sized array of bits implementing
 CReshapeTile
 CReshapeTile< Tile_, kAccessSize_, true >
 CShapeA Shape implementing Layout Concept describing the dimensions of a cube
 CShapeAdd
 CShapeCountCompute derived counted of a Layout Concept based class
 CShapeDiv
 CShapeMax
 CShapeMin
 CShapeMul
 CShapeScale
 CShapeStrides
 CShapeSub
 Csqrt_est
 CStorageType
 CStorageType< 1 >
 CStorageType< 2 >
 CStorageType< 4 >
 CStore
 CStore< double, 2, Memory_, true, 16 >
 CStore< Scalar_, Lanes_, Memory_, true, 16 >
 CStore< Scalar_, Lanes_, Memory_, true, 4 >
 CStore< Scalar_, Lanes_, Memory_, true, 8 >
 CTensorRefStructure modeling a pointer and stride into a tensor
 CTensorViewHost-side reference implementation of tensor operations
 CTiledThreadOffsetBasic thread offset function computed from a thread shape
 ▶CTileIteratorBaseIterator for accessing a stripmined tile in memory
 ▶CTileLoadIteratorAn iterator implementing Tile Load Iterator Concept for loading a tile from memory
 ▶CTileStoreIteratorAn iterator implementing Tile Store Iterator Concept for storing a tile to memory
 CTileTraitsA template defining Tile Traits Concept
 CTileTraitsContiguousMajor
 CTileTraitsStandardChooses 'best' shape to enable warp raking along contiguous dimension if possible
 CTileTraitsStrideMajor
 ▶CTileTraitsWarpRakeTiling in which warps rake across the contiguous dimension
 CTrivialPredicateTileAdapterAlways returns true predicate
 CVector
 CVector< half, kLanes_ >
 CVectorize
 CVectorize< Element_, 1 >
 CVectorTraitsTraits describing properties of vectors and scalar-as-vectors
 CVectorTraits< Vector< T, Lanes > >Partial specialization for actual cutlass::Vector
 CVectorTraits< Vector< T, Lanes > const >Partial specialization for actual cutlass::Vector
 ▶Ncutlass
 CDebugType
 CDebugValue
-Generated on Thu May 3 2018 16:36:15 for Cutlass by   +Generated on Wed Sep 19 2018 13:58:53 for Cutlass by   doxygen 1.8.14
diff --git a/docs/classcutlass_1_1PredicateVector_1_1ConstIterator-members.html b/docs/classcutlass_1_1PredicateVector_1_1ConstIterator-members.html index 860cd05c..18f59fc0 100644 --- a/docs/classcutlass_1_1PredicateVector_1_1ConstIterator-members.html +++ b/docs/classcutlass_1_1PredicateVector_1_1ConstIterator-members.html @@ -91,7 +91,7 @@ $(function() {
-Generated on Thu May 3 2018 16:36:13 for Cutlass by   +Generated on Wed Sep 19 2018 13:58:51 for Cutlass by   doxygen 1.8.14
diff --git a/docs/classcutlass_1_1PredicateVector_1_1ConstIterator.html b/docs/classcutlass_1_1PredicateVector_1_1ConstIterator.html index 1fbdc759..7e7089a0 100644 --- a/docs/classcutlass_1_1PredicateVector_1_1ConstIterator.html +++ b/docs/classcutlass_1_1PredicateVector_1_1ConstIterator.html @@ -381,7 +381,7 @@ template<int kPredicates_, int kPredicatesPerByte_ = 4, int kPredicateStart_
-Generated on Thu May 3 2018 16:36:13 for Cutlass by   +Generated on Wed Sep 19 2018 13:58:51 for Cutlass by   doxygen 1.8.14
diff --git a/docs/classcutlass_1_1PredicateVector_1_1Iterator-members.html b/docs/classcutlass_1_1PredicateVector_1_1Iterator-members.html index ca3ff04a..73d0ebca 100644 --- a/docs/classcutlass_1_1PredicateVector_1_1Iterator-members.html +++ b/docs/classcutlass_1_1PredicateVector_1_1Iterator-members.html @@ -93,7 +93,7 @@ $(function() {
-Generated on Thu May 3 2018 16:36:13 for Cutlass by   +Generated on Wed Sep 19 2018 13:58:51 for Cutlass by   doxygen 1.8.14
diff --git a/docs/classcutlass_1_1PredicateVector_1_1Iterator.html b/docs/classcutlass_1_1PredicateVector_1_1Iterator.html index 42a06938..2cbc797d 100644 --- a/docs/classcutlass_1_1PredicateVector_1_1Iterator.html +++ b/docs/classcutlass_1_1PredicateVector_1_1Iterator.html @@ -443,7 +443,7 @@ template<int kPredicates_, int kPredicatesPerByte_ = 4, int kPredicateStart_
-Generated on Thu May 3 2018 16:36:13 for Cutlass by   +Generated on Wed Sep 19 2018 13:58:51 for Cutlass by   doxygen 1.8.14
diff --git a/docs/classcutlass_1_1TensorRef-members.html b/docs/classcutlass_1_1TensorRef-members.html index 4bf37ad1..202c9ab4 100644 --- a/docs/classcutlass_1_1TensorRef-members.html +++ b/docs/classcutlass_1_1TensorRef-members.html @@ -73,35 +73,52 @@ $(function() {
-
cutlass::TensorRef< Storage_, Rank_ > Member List
+
cutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ > Member List
-

This is the complete list of members for cutlass::TensorRef< Storage_, Rank_ >, including all inherited members.

+

This is the complete list of members for cutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >, including all inherited members.

- - - - - - - - - - - - - - - - - - - + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + +
advance(Coord< Rank > const &b)cutlass::TensorRef< Storage_, Rank_ >inline
at(Coord< Rank > const &coord) constcutlass::TensorRef< Storage_, Rank_ >inline
at(int idx) constcutlass::TensorRef< Storage_, Rank_ >inline
convert()cutlass::TensorRef< Storage_, Rank_ >inline
data() constcutlass::TensorRef< Storage_, Rank_ >inline
good() constcutlass::TensorRef< Storage_, Rank_ >inline
leading_dim() constcutlass::TensorRef< Storage_, Rank_ >inline
offset(Coord< Rank > const &coord) constcutlass::TensorRef< Storage_, Rank_ >inline
operator+(Coord< Rank > const &b) constcutlass::TensorRef< Storage_, Rank_ >inline
operator-(Coord< Rank > const &b) constcutlass::TensorRef< Storage_, Rank_ >inline
operator[](Coord< Rank > const &coord) constcutlass::TensorRef< Storage_, Rank_ >inline
operator[](int idx) constcutlass::TensorRef< Storage_, Rank_ >inline
Rankcutlass::TensorRef< Storage_, Rank_ >static
reset(Storage *ptr=nullptr, Coord< Rank > stride=Coord< Rank >(0))cutlass::TensorRef< Storage_, Rank_ >inline
Storage typedefcutlass::TensorRef< Storage_, Rank_ >
stride() constcutlass::TensorRef< Storage_, Rank_ >inline
stride(int dim) constcutlass::TensorRef< Storage_, Rank_ >inline
TensorRef()cutlass::TensorRef< Storage_, Rank_ >inline
TensorRef(Storage *ptr, Coord< Rank > stride)cutlass::TensorRef< Storage_, Rank_ >inline
add_pointer_offset(LongIndex delta)cutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >inline
at(TensorCoord const &coord) constcutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >inline
at(LongIndex idx) constcutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >inline
const_ref() constcutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >inline
ConstTensorRef typedefcutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >
Coord_t typedefcutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >
data() constcutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >inline
good() constcutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >inline
Index typedefcutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >
kRankcutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >static
kStorageRankcutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >static
leading_dim(int idx=0) constcutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >inline
LongIndex typedefcutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >
map(TensorCoord const &coord) constcutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >inline
MapFunc typedefcutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >
offset(TensorCoord const &coord) constcutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >inline
operator+(TensorCoord const &b) constcutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >inline
operator+=(TensorCoord const &b)cutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >inline
operator-(TensorCoord const &b) constcutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >inline
operator-=(TensorCoord const &b)cutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >inline
operator[](TensorCoord const &coord) constcutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >inline
operator[](LongIndex idx) constcutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >inline
Rankcutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >static
reset(Storage *ptr=nullptr)cutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >inline
reset(Storage *ptr, StorageCoord const &stride)cutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >inline
Storage typedefcutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >
StorageCoord typedefcutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >
stride() constcutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >inline
stride(int dim) constcutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >inline
StrideVector typedefcutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >
TensorCoord typedefcutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >
TensorRef(Storage *ptr=nullptr)cutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >inline
TensorRef(Storage *ptr, Index ldm)cutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >inline
TensorRef(Storage *ptr, StrideVector const &stride)cutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >inline
TensorRef(Storage *ptr, StorageCoord const &stride)cutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >inline
TensorRef(TensorRef< typename platform::remove_const< Storage >::type, kRank, MapFunc, kStorageRank, Index, LongIndex > const &ref)cutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >inline
-Generated on Thu May 3 2018 16:36:13 for Cutlass by   +Generated on Wed Sep 19 2018 13:58:51 for Cutlass by   doxygen 1.8.14
diff --git a/docs/classcutlass_1_1TensorRef.html b/docs/classcutlass_1_1TensorRef.html index 05a9b3dd..1053ca0a 100644 --- a/docs/classcutlass_1_1TensorRef.html +++ b/docs/classcutlass_1_1TensorRef.html @@ -5,7 +5,7 @@ -Cutlass: cutlass::TensorRef< Storage_, Rank_ > Class Template Reference +Cutlass: cutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ > Class Template Reference @@ -78,93 +78,278 @@ $(function() { Static Public Attributes | List of all members
-
cutlass::TensorRef< Storage_, Rank_ > Class Template Reference
+
cutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ > Class Template Reference
-

Structure modeling a pointer and stride into a tensor. -

-

#include <tensor_ref.h>

+
+Inheritance diagram for cutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >:
+
+
+ + +cutlass::TensorRefBatchStrided< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ > +cutlass::TensorView< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ > + +
- - - + + + + + + + + + + + + + + + + + + + + + + + + + +

Public Types

typedef Storage_ Storage
 Data type of individual access. More...
 
typedef Storage_ Storage
 Data type of individual access. More...
 
typedef MapFunc_ MapFunc
 Mapping function from logical coordinate to internal n-D array. More...
 
typedef Index_ Index
 Index type. More...
 
typedef LongIndex_ LongIndex
 Typically, strides in memory can be very large. More...
 
typedef Coord< kRank > TensorCoord
 Coordinate in logical tensor space. More...
 
typedef Coord< kStorageRank > StorageCoord
 Coordinate in storage n-D array. More...
 
typedef Coord< kStorageRank - 1 > StrideVector
 
typedef TensorRef< typename platform::remove_const< Storage >::type const, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ > ConstTensorRef
 Tensor reference to of constant value. More...
 
typedef TensorCoord Coord_t
 Coordinate in logical tensor space. More...
 
- - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + +

Public Member Functions

CUTLASS_HOST_DEVICE TensorRef ()
 Default ctor. More...
 
CUTLASS_HOST_DEVICE TensorRef (Storage *ptr, Coord< Rank > stride)
 Constructs from a pointer, size, and stride. More...
 
CUTLASS_HOST_DEVICE void reset (Storage *ptr=nullptr, Coord< Rank > stride=Coord< Rank >(0))
 Updates the pointer, stride, and location within a TensorRef. More...
 
template<typename T >
TensorRef< T, Rank > convert ()
 Conversion function. More...
 
CUTLASS_HOST_DEVICE bool good () const
 Returns true if the TensorRef may be safely accessed. More...
 
CUTLASS_HOST_DEVICE Storage * data () const
 Returns the pointer to referenced data. More...
 
CUTLASS_HOST_DEVICE Coord< Rank > const & stride () const
 Returns the stride of the tensor. More...
 
CUTLASS_HOST_DEVICE int const & stride (int dim) const
 Returns the stride of the tensor in the given dimension. More...
 
CUTLASS_HOST_DEVICE int leading_dim () const
 Returns the maximum stride element as the 'leading dimension'. More...
 
CUTLASS_HOST_DEVICE long long offset (Coord< Rank > const &coord) const
 Computes the offset of an index from the origin of the tensor. More...
 
CUTLASS_HOST_DEVICE Storage & at (Coord< Rank > const &coord) const
 Returns a reference to the element at a given Coord. More...
 
Storage & operator[] (Coord< Rank > const &coord) const
 Element-wise accessor. More...
 
CUTLASS_HOST_DEVICE Storage & at (int idx) const
 Returns a reference to the element at a given Coord. More...
 
Storage & operator[] (int idx) const
 Element-wise accessor. More...
 
CUTLASS_HOST_DEVICE TensorRef & advance (Coord< Rank > const &b)
 Adds an offset to the pointer. More...
 
CUTLASS_HOST_DEVICE TensorRef operator+ (Coord< Rank > const &b) const
 Returns a TensorRef offset by a given amount. More...
 
CUTLASS_HOST_DEVICE TensorRef operator- (Coord< Rank > const &b) const
 Returns a TensorRef offset by a given amount. More...
 
CUTLASS_HOST_DEVICE TensorRef (Storage *ptr=nullptr)
 Helper for 1-D memory. All higher ranks are projected onto the fastest changing rank. More...
 
CUTLASS_HOST_DEVICE TensorRef (Storage *ptr, Index ldm)
 Helper to construct from a pointer and single stride element for 2-D pitch linear memory. More...
 
CUTLASS_HOST_DEVICE TensorRef (Storage *ptr, StrideVector const &stride)
 Constructs from a single pointer and stride vector. More...
 
CUTLASS_HOST_DEVICE TensorRef (Storage *ptr, StorageCoord const &stride)
 
CUTLASS_HOST_DEVICE TensorRef (TensorRef< typename platform::remove_const< Storage >::type, kRank, MapFunc, kStorageRank, Index, LongIndex > const &ref)
 Enables conversion from TensorRef of non-const type. More...
 
CUTLASS_HOST_DEVICE ConstTensorRef const_ref () const
 Returns a reference to constant-valued tensor. More...
 
CUTLASS_HOST_DEVICE void reset (Storage *ptr=nullptr)
 Updates only the pointer. More...
 
CUTLASS_HOST_DEVICE void reset (Storage *ptr, StorageCoord const &stride)
 Updates the pointer, stride, and location within a TensorRef. More...
 
CUTLASS_HOST_DEVICE bool good () const
 Returns true if the TensorRef may be safely accessed. More...
 
CUTLASS_HOST_DEVICE Storage * data () const
 Returns the pointer to referenced data. More...
 
CUTLASS_HOST_DEVICE StorageCoord stride () const
 Returns the stride of the tensor. More...
 
CUTLASS_HOST_DEVICE Index stride (int dim) const
 Returns the stride of the tensor in the given dimension. More...
 
CUTLASS_HOST_DEVICE Index leading_dim (int idx=0) const
 Returns the maximum stride element as the 'leading dimension'. More...
 
CUTLASS_HOST_DEVICE StorageCoord map (TensorCoord const &coord) const
 Maps a logical coordinate to an n-D array in memory. More...
 
CUTLASS_HOST_DEVICE LongIndex offset (TensorCoord const &coord) const
 Computes the offset of an index from the origin of the tensor. More...
 
CUTLASS_HOST_DEVICE Storage & at (TensorCoord const &coord) const
 Returns a reference to the element at a given Coord. More...
 
CUTLASS_HOST_DEVICE Storage & at (LongIndex idx) const
 Returns a reference to the element at a given linear index. More...
 
CUTLASS_HOST_DEVICE Storage & operator[] (TensorCoord const &coord) const
 Returns a reference to the element at a given Coord. More...
 
CUTLASS_HOST_DEVICE Storage & operator[] (LongIndex idx) const
 Returns a reference to the element at a given linear index. More...
 
CUTLASS_HOST_DEVICE TensorRef & add_pointer_offset (LongIndex delta)
 Adds an offset to each pointer. More...
 
CUTLASS_HOST_DEVICE TensorRef operator+ (TensorCoord const &b) const
 Returns a TensorRef offset by a given amount. More...
 
CUTLASS_HOST_DEVICE TensorRef & operator+= (TensorCoord const &b)
 Returns a TensorRef offset by a given amount. More...
 
CUTLASS_HOST_DEVICE TensorRef operator- (TensorCoord const &b) const
 Returns a TensorRef offset by a given amount. More...
 
CUTLASS_HOST_DEVICE TensorRef & operator-= (TensorCoord const &b)
 Returns a TensorRef offset by a given amount. More...
 
- - - + + + + + + + + +

Static Public Attributes

static int const Rank = Rank_
 Rank of tensor. More...
 
static int const kRank = Rank_
 Logical rank of tensor index space. More...
 
static int const kStorageRank = StorageRank_
 Rank of internal storage. More...
 
static int const Rank = kRank
 Logical rank of tensor index space. More...
 

Member Typedef Documentation

- -

◆ Storage

+ +

◆ ConstTensorRef

-template<typename Storage_, int Rank_>
+template<typename Storage_, int Rank_, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>
- + + +
typedef Storage_ cutlass::TensorRef< Storage_, Rank_ >::Storagetypedef TensorRef< typename platform::remove_const<Storage>::type const, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_> cutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::ConstTensorRef
+
+ +
+
+ +

◆ Coord_t

+ +
+
+
+template<typename Storage_, int Rank_, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>
+ + + + +
typedef TensorCoord cutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::Coord_t
+
+

Require at least rank=1. Mathematically, a rank=0 tensor would be considered to be a scalar, but degenerate cases such as these are difficult to accommodate without extensive C++ metaprogramming or support for zero-length arrays.

+ +
+
+ +

◆ Index

+ +
+
+
+template<typename Storage_, int Rank_, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>
+ + + + +
typedef Index_ cutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::Index
+
+ +
+
+ +

◆ LongIndex

+ +
+
+
+template<typename Storage_, int Rank_, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>
+ + + + +
typedef LongIndex_ cutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::LongIndex
+
+ +
+
+ +

◆ MapFunc

+ +
+
+
+template<typename Storage_, int Rank_, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>
+ + + + +
typedef MapFunc_ cutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::MapFunc
+
+ +
+
+ +

◆ Storage

+ +
+
+
+template<typename Storage_, int Rank_, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>
+ + + + +
typedef Storage_ cutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::Storage
+
+ +
+
+ +

◆ StorageCoord

+ +
+
+
+template<typename Storage_, int Rank_, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>
+ + + + +
typedef Coord<kStorageRank> cutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::StorageCoord
+
+ +
+
+ +

◆ StrideVector

+ +
+
+
+template<typename Storage_, int Rank_, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>
+ + + + +
typedef Coord<kStorageRank - 1> cutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::StrideVector
+
+

Stride vector in storage coordinage space - assumes least significant stride is 1 and does not store it.

+ +
+
+ +

◆ TensorCoord

+ +
+
+
+template<typename Storage_, int Rank_, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>
+ + +
typedef Coord<kRank> cutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::TensorCoord
@@ -172,21 +357,22 @@ template<typename Storage_, int Rank_>

Constructor & Destructor Documentation

- -

◆ TensorRef() [1/2]

+ +

◆ TensorRef() [1/5]

-template<typename Storage_, int Rank_>
+template<typename Storage_, int Rank_, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>
-Generated on Thu May 3 2018 16:36:12 for Cutlass by   +Generated on Wed Sep 19 2018 13:58:50 for Cutlass by   doxygen 1.8.14
diff --git a/docs/hgemm__global__tile_8h.html b/docs/hgemm__global__tile_8h.html index b62b8c14..0b2e247f 100644 --- a/docs/hgemm__global__tile_8h.html +++ b/docs/hgemm__global__tile_8h.html @@ -82,10 +82,10 @@ $(function() {

Tile traits used to construct global tile iterator for HGEMM. This is intended to partition the thread block-level tile into 2D subtiles loaded by the threads and facilitate memory accesses larger than 16 bits. More...

-
#include <cutlass/coord.h>
-#include <cutlass/gemm/gemm_global_tile.h>
-#include <cutlass/matrix_traits.h>
-#include <cutlass/reshape_tile.h>
+
#include "cutlass/coord.h"
+#include "cutlass/gemm/gemm_global_tile.h"
+#include "cutlass/matrix_traits.h"
+#include "cutlass/reshape_tile.h"

Go to the source code of this file.

- + - + +
CUTLASS_HOST_DEVICE cutlass::TensorRef< Storage_, Rank_ >::TensorRef CUTLASS_HOST_DEVICE cutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::TensorRef ()Storage * ptr = nullptr)
@@ -199,27 +385,65 @@ template<typename Storage_, int Rank_> - -

◆ TensorRef() [2/2]

+ +

◆ TensorRef() [2/5]

-template<typename Storage_, int Rank_>
+template<typename Storage_, int Rank_, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>
+ + +
- + - + - + + + + + + + + +
CUTLASS_HOST_DEVICE cutlass::TensorRef< Storage_, Rank_ >::TensorRef CUTLASS_HOST_DEVICE cutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::TensorRef (Storage * Storage *  ptr,
Coord< Rank > Index ldm 
)
+ 		+inline
+
+ +
+ + +

◆ TensorRef() [3/5]

+ +
+
+
+template<typename Storage_, int Rank_, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>
+ + +
-Generated on Thu May 3 2018 16:36:12 for Cutlass by   +Generated on Wed Sep 19 2018 13:58:50 for Cutlass by   doxygen 1.8.14
diff --git a/docs/group__tile__traits__concept.html b/docs/group__tile__traits__concept.html index 16e4bd8a..6c051696 100644 --- a/docs/group__tile__traits__concept.html +++ b/docs/group__tile__traits__concept.html @@ -77,7 +77,7 @@ $(function() {
+ + + + + + + + + + + @@ -237,363 +461,541 @@ template<typename Storage_, int Rank_> -

Member Function Documentation

- -

◆ advance()

+ +

◆ TensorRef() [4/5]

-template<typename Storage_, int Rank_>
+template<typename Storage_, int Rank_, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>
CUTLASS_HOST_DEVICE cutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::TensorRef (Storage * ptr,
StrideVector const &  stride 
- - -
- + - - - - -
CUTLASS_HOST_DEVICE TensorRef& cutlass::TensorRef< Storage_, Rank_ >::advance CUTLASS_HOST_DEVICE cutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::TensorRef (Coord< Rank > const & b)
- 		-inline
-
- -
- - -

◆ at() [1/2]

- -
-
-
-template<typename Storage_, int Rank_>
- - - - - -
- - - - - - - - -
CUTLASS_HOST_DEVICE Storage& cutlass::TensorRef< Storage_, Rank_ >::at (Coord< Rank > const & coord) const
- 		-inline
-
- -
-
- -

◆ at() [2/2]

- -
-
-
-template<typename Storage_, int Rank_>
- - - - - -
- - - - - - - - -
CUTLASS_HOST_DEVICE Storage& cutlass::TensorRef< Storage_, Rank_ >::at (int idx) const
- 		-inline
-
- -
-
- -

◆ convert()

- -
-
-
-template<typename Storage_, int Rank_>
-
-template<typename T >
- - - - - -
- - - - - - - -
TensorRef<T, Rank> cutlass::TensorRef< Storage_, Rank_ >::convert ()
- 		-inline
-
- -
-
- -

◆ data()

- -
-
-
-template<typename Storage_, int Rank_>
- - - - - -
- - - - - - - -
CUTLASS_HOST_DEVICE Storage* cutlass::TensorRef< Storage_, Rank_ >::data () const
- 		-inline
-
- -
-
- -

◆ good()

- -
-
-
-template<typename Storage_, int Rank_>
- - - - - -
- - - - - - - -
CUTLASS_HOST_DEVICE bool cutlass::TensorRef< Storage_, Rank_ >::good () const
- 		-inline
-
- -
-
- -

◆ leading_dim()

- -
-
-
-template<typename Storage_, int Rank_>
- - - - - -
- - - - - - - -
CUTLASS_HOST_DEVICE int cutlass::TensorRef< Storage_, Rank_ >::leading_dim () const
- 		-inline
-
- -
-
- -

◆ offset()

- -
-
-
-template<typename Storage_, int Rank_>
- - - - - -
- - - - - - - - -
CUTLASS_HOST_DEVICE long long cutlass::TensorRef< Storage_, Rank_ >::offset (Coord< Rank > const & coord) const
- 		-inline
-
- -
-
- -

◆ operator+()

- -
-
-
-template<typename Storage_, int Rank_>
- - - - - -
- - - - - - - - -
CUTLASS_HOST_DEVICE TensorRef cutlass::TensorRef< Storage_, Rank_ >::operator+ (Coord< Rank > const & b) const
- 		-inline
-
- -
-
- -

◆ operator-()

- -
-
-
-template<typename Storage_, int Rank_>
- - - - - -
- - - - - - - - -
CUTLASS_HOST_DEVICE TensorRef cutlass::TensorRef< Storage_, Rank_ >::operator- (Coord< Rank > const & b) const
- 		-inline
-
- -
-
- -

◆ operator[]() [1/2]

- -
-
-
-template<typename Storage_, int Rank_>
- - - - - -
- - - - - - - - -
Storage& cutlass::TensorRef< Storage_, Rank_ >::operator[] (Coord< Rank > const & coord) const
- 		-inline
-
- -
-
- -

◆ operator[]() [2/2]

- -
-
-
-template<typename Storage_, int Rank_>
- - - - - -
- - - - - - - - -
Storage& cutlass::TensorRef< Storage_, Rank_ >::operator[] (int idx) const
- 		-inline
-
- -
-
- -

◆ reset()

- -
-
-
-template<typename Storage_, int Rank_>
- - - + + +
- - - - - - + + - - + + + + + + + + +
CUTLASS_HOST_DEVICE void cutlass::TensorRef< Storage_, Rank_ >::reset (Storage * ptr = nullptr, Storage * ptr,
Coord< Rank > stride = Coord<Rank>(0) StorageCoord const & stride 
)
+ 		+inline
+
+

Constructs from a pointer and a stride vector of size kRank. If fastest changing stride is not 1, construction fails and subsequent calls to good() will return false.

+ +
+
+ +

◆ TensorRef() [5/5]

+ +
+
+
+template<typename Storage_, int Rank_, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>
+ + + + + +
+ + + + + + + + +
CUTLASS_HOST_DEVICE cutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::TensorRef (TensorRef< typename platform::remove_const< Storage >::type, kRank, MapFunc, kStorageRank, Index, LongIndex > const & ref)
+ 		+inline
+
+ +
+
+

Member Function Documentation

+ +

◆ add_pointer_offset()

+ +
+
+
+template<typename Storage_, int Rank_, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>
+ + + + + +
+ + + + + + + + +
CUTLASS_HOST_DEVICE TensorRef& cutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::add_pointer_offset (LongIndex delta)
+ 		+inline
+
+ +
+
+ +

◆ at() [1/2]

+ +
+
+
+template<typename Storage_, int Rank_, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>
+ + + + + +
+ + + + + + + + +
CUTLASS_HOST_DEVICE Storage& cutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::at (TensorCoord const & coord) const
+ 		+inline
+
+ +
+
+ +

◆ at() [2/2]

+ +
+
+
+template<typename Storage_, int Rank_, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>
+ + + + + +
+ + + + + + + + +
CUTLASS_HOST_DEVICE Storage& cutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::at (LongIndex idx) const
+ 		+inline
+
+ +
+
+ +

◆ const_ref()

+ +
+
+
+template<typename Storage_, int Rank_, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>
+ + + + + +
+ + + + + + + +
CUTLASS_HOST_DEVICE ConstTensorRef cutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::const_ref () const
+ 		+inline
+
+ +
+
+ +

◆ data()

+ +
+
+
+template<typename Storage_, int Rank_, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>
+ + + + + +
+ + + + + + + +
CUTLASS_HOST_DEVICE Storage* cutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::data () const
+ 		+inline
+
+ +
+
+ +

◆ good()

+ +
+
+
+template<typename Storage_, int Rank_, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>
+ + + + + +
+ + + + + + + +
CUTLASS_HOST_DEVICE bool cutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::good () const
+ 		+inline
+
+ +
+
+ +

◆ leading_dim()

+ +
+
+
+template<typename Storage_, int Rank_, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>
+ + + + + +
+ + + + + + + + +
CUTLASS_HOST_DEVICE Index cutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::leading_dim (int idx = 0) const
+ 		+inline
+
+ +
+
+ +

◆ map()

+ +
+
+
+template<typename Storage_, int Rank_, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>
+ + + + + +
+ + + + + + + + +
CUTLASS_HOST_DEVICE StorageCoord cutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::map (TensorCoord const & coord) const
+ 		+inline
+
+ +
+
+ +

◆ offset()

+ +
+
+
+template<typename Storage_, int Rank_, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>
+ + + + + +
+ + + + + + + + +
CUTLASS_HOST_DEVICE LongIndex cutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::offset (TensorCoord const & coord) const
+ 		+inline
+
+ +
+
+ +

◆ operator+()

+ +
+
+
+template<typename Storage_, int Rank_, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>
+ + + + + +
+ + + + + + + + +
CUTLASS_HOST_DEVICE TensorRef cutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::operator+ (TensorCoord const & b) const
+ 		+inline
+
+ +
+
+ +

◆ operator+=()

+ +
+
+
+template<typename Storage_, int Rank_, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>
+ + + + + +
+ + + + + + + + +
CUTLASS_HOST_DEVICE TensorRef& cutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::operator+= (TensorCoord const & b)
+ 		+inline
+
+ +
+
+ +

◆ operator-()

+ +
+
+
+template<typename Storage_, int Rank_, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>
+ + + + + +
+ + + + + + + + +
CUTLASS_HOST_DEVICE TensorRef cutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::operator- (TensorCoord const & b) const
+ 		+inline
+
+ +
+
+ +

◆ operator-=()

+ +
+
+
+template<typename Storage_, int Rank_, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>
+ + + + + +
+ + + + + + + + +
CUTLASS_HOST_DEVICE TensorRef& cutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::operator-= (TensorCoord const & b)
+ 		+inline
+
+ +
+
+ +

◆ operator[]() [1/2]

+ +
+
+
+template<typename Storage_, int Rank_, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>
+ + + + + +
+ + + + + + + + +
CUTLASS_HOST_DEVICE Storage& cutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::operator[] (TensorCoord const & coord) const
+ 		+inline
+
+ +
+
+ +

◆ operator[]() [2/2]

+ +
+
+
+template<typename Storage_, int Rank_, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>
+ + + + + +
+ + + + + + + + +
CUTLASS_HOST_DEVICE Storage& cutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::operator[] (LongIndex idx) const
+ 		+inline
+
+ +
+
+ +

◆ reset() [1/2]

+ +
+
+
+template<typename Storage_, int Rank_, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>
+ + + + + +
+ + + + + + + + +
CUTLASS_HOST_DEVICE void cutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::reset (Storage * ptr = nullptr)
+ 		+inline
+
+ +
+
+ +

◆ reset() [2/2]

+ +
+
+
+template<typename Storage_, int Rank_, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>
+ + +
-Generated on Thu May 3 2018 16:36:12 for Cutlass by   +Generated on Wed Sep 19 2018 13:58:50 for Cutlass by   doxygen 1.8.14
diff --git a/docs/group__tile__store__iterator__concept.html b/docs/group__tile__store__iterator__concept.html index bde54053..992a7ca3 100644 --- a/docs/group__tile__store__iterator__concept.html +++ b/docs/group__tile__store__iterator__concept.html @@ -77,7 +77,7 @@ $(function() {
+ + + + + + + + + + + + @@ -610,19 +1012,19 @@ template<typename Storage_, int Rank_> - -

◆ stride() [1/2]

+ +

◆ stride() [1/2]

-template<typename Storage_, int Rank_>
+template<typename Storage_, int Rank_, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>
CUTLASS_HOST_DEVICE void cutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::reset (Storage * ptr,
StorageCoord const & stride 
-Generated on Thu May 3 2018 16:36:12 for Cutlass by   +Generated on Wed Sep 19 2018 13:58:50 for Cutlass by   doxygen 1.8.14
diff --git a/docs/group__fragment__iterator__concept.html b/docs/group__fragment__iterator__concept.html index dc89e72e..e08d36e4 100644 --- a/docs/group__fragment__iterator__concept.html +++ b/docs/group__fragment__iterator__concept.html @@ -91,7 +91,7 @@ Classes
-Generated on Thu May 3 2018 16:36:12 for Cutlass by   +Generated on Wed Sep 19 2018 13:58:50 for Cutlass by   doxygen 1.8.14
diff --git a/docs/group__layout__concept.html b/docs/group__layout__concept.html index 3fe8532c..66a82881 100644 --- a/docs/group__layout__concept.html +++ b/docs/group__layout__concept.html @@ -100,7 +100,7 @@ Classes
-Generated on Thu May 3 2018 16:36:12 for Cutlass by   +Generated on Wed Sep 19 2018 13:58:50 for Cutlass by   doxygen 1.8.14
diff --git a/docs/group__predicate__iterator__concept.html b/docs/group__predicate__iterator__concept.html index 95c1ef2e..9c3b7108 100644 --- a/docs/group__predicate__iterator__concept.html +++ b/docs/group__predicate__iterator__concept.html @@ -98,7 +98,7 @@ Classes
-Generated on Thu May 3 2018 16:36:12 for Cutlass by   +Generated on Wed Sep 19 2018 13:58:50 for Cutlass by   doxygen 1.8.14
diff --git a/docs/group__predicate__tile__adapter.html b/docs/group__predicate__tile__adapter.html index a4b80992..8ab28fed 100644 --- a/docs/group__predicate__tile__adapter.html +++ b/docs/group__predicate__tile__adapter.html @@ -80,7 +80,7 @@ $(function() {
-Generated on Thu May 3 2018 16:36:12 for Cutlass by   +Generated on Wed Sep 19 2018 13:58:50 for Cutlass by   doxygen 1.8.14
diff --git a/docs/group__predicate__vector__concept.html b/docs/group__predicate__vector__concept.html index 5147870e..cf4fd5b2 100644 --- a/docs/group__predicate__vector__concept.html +++ b/docs/group__predicate__vector__concept.html @@ -92,7 +92,7 @@ Classes
-Generated on Thu May 3 2018 16:36:12 for Cutlass by   +Generated on Wed Sep 19 2018 13:58:50 for Cutlass by   doxygen 1.8.14
diff --git a/docs/group__tile__load__iterator__concept.html b/docs/group__tile__load__iterator__concept.html index 2bc4b4e3..edc49281 100644 --- a/docs/group__tile__load__iterator__concept.html +++ b/docs/group__tile__load__iterator__concept.html @@ -77,7 +77,7 @@ $(function() {
- + @@ -637,19 +1039,19 @@ template<typename Storage_, int Rank_> - -

◆ stride() [2/2]

+ +

◆ stride() [2/2]

-template<typename Storage_, int Rank_>
+template<typename Storage_, int Rank_, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>
CUTLASS_HOST_DEVICE Coord<Rank> const& cutlass::TensorRef< Storage_, Rank_ >::stride CUTLASS_HOST_DEVICE StorageCoord cutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::stride ( ) const
+ + +
- + @@ -666,19 +1068,67 @@ template<typename Storage_, int Rank_>

Member Data Documentation

- -

◆ Rank

+ +

◆ kRank

-template<typename Storage_, int Rank_>
+template<typename Storage_, int Rank_, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>
CUTLASS_HOST_DEVICE int const& cutlass::TensorRef< Storage_, Rank_ >::stride CUTLASS_HOST_DEVICE Index cutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::stride ( int  dim)
+ + +
- + + +
int const cutlass::TensorRef< Storage_, Rank_ >::Rank = Rank_int const cutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::kRank = Rank_
+ 		+static
+
+ +
+ + +

◆ kStorageRank

+ +
+
+
+template<typename Storage_, int Rank_, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>
+ + + + + +
+ + + + +
int const cutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::kStorageRank = StorageRank_
+ 		+static
+
+ +
+
+ +

◆ Rank

+ +
+
+
+template<typename Storage_, int Rank_, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>
+ + + @@ -696,7 +1146,7 @@ template<typename Storage_, int Rank_>
-Generated on Thu May 3 2018 16:36:13 for Cutlass by   +Generated on Wed Sep 19 2018 13:58:51 for Cutlass by   doxygen 1.8.14
diff --git a/docs/classcutlass_1_1TensorRef.png b/docs/classcutlass_1_1TensorRef.png new file mode 100644 index 0000000000000000000000000000000000000000..f8caaa61d05cb4460b85cc8dce3844d1f38cc22d GIT binary patch literal 2451 zcmb_edo+~o7N^o>+*VO2bjr{umz5A@NMF)m7=}T}=a$Aa?lco4A7OnXqb7+_#&sk{ zavKrm(;1UnVvcbcV?;F}_uMaYUh1r~*7@VC^T&DCdf#Wi&$IVl``N$!+t2$a!q!q6 zbQB~eCMFHLcEv$VY!3yH5&L%o@0-ap5x}H~uyM2yi9~=nyZDXP9Ptp4b{-~^>GB$f z0zUVLJJ=${0Hy8G+vcbzCU)Q??25S~PJD|0DKh^;<^4}d-t-}n`n^h@cy*?XVkkM2 zQ@MEa*sbd84#htdKr2Kq#YAt##9YNhFPs7-Csq;yHy1wzQ8E{Iwx`7HvPI~{?JCRD zmymdt~SZ^be->4F`xap$N4|A0RECjyrsB`!BLQg&=(7J+LIW>JuV!d@~=jcbK zyve$liEyv@n=iY~j%3qdkt_%pz1|WA2 za*ML9`XfnPO8Td;imo!g)W;gNUe%XnZcKInc1GYnANnQ8L&i7md71%mu#ij0&VfwF z#=E$;SQ)`ohwi$Kt^)pUa+Zk?e=@sdHLgY**>y_p$WZDy3TL()v-jKJ9|5UOeL5qd zz1`JlX=?aTCYlrFl~thu896O;9VtYlDl9x2Do=c}(7|4cL<@A6zcM57v@YzcYARNx zJ8GhO-z7UAc>AmqftMIuHR#ieePE_^JiN399g*<8J^cN%HcC!xi4g|Q>GsHK+w2rB zPw$a(Im@QoUAa@Biau|*>b*h#-DQ5YA*;Lj{?c{Mbbzcg1DzqHkNeR!MAXjXyj#zzR(-jK)e31n+vWLY#N&I(-~Z z+uh>G1+ZnCW9r(Y&~W}j55=(AkX~%8NmqMo@_pGVTeaIiNe;^kerQ^Ad}8w8J}s9b zBLp8s;;!FuL7vN3$&^t-ZUd{5RMwJpL$uLwU>(`Fz$p-Z9qJ_9oc5w);hu(F>(oAE z6V3VDME*?ZRga0NkshbPe#;cbY-7%PtVeZ;feB*juL+98voI|wKT>vyOX)_lKGi=i zeh0NV1}4BmbfPi=Ue+2{rs6i0oZl2y(!QOyl`Mlr{A8Cb<6gZ+m0F}FMNLhoB*9;~ zy$ZDtR?p4!;kR0ZsQ2dcvI81tQ{%jUn^TiNW7dp;*z)-9Eb*(o8t5|wyk2sfiNzKOKj}= zr7E#~UJ3bG$!%CckP!t00FOL^shb0^0CXY zL)8oP8&z0E((dhMRgysvr25shP;?Frk=zS2Jo7TZc$OQt>ds16&M?971|50}y4L>2 zZ2U3Uv*Lwt5jFgz>6{77(8ysbcYe}vL$%9){7Qr)| zB>q!j`sgF!lK+@fmT%|_x9^UuybP=$fXu(Jq*4=9+?Y)xri`F&xmnl^ue+ocsZEzI zAtLx!P1130T4Am2DsLlj3YYztxjlJz#o@0|+^V7YnC_XV28v)5 z^t_EXRJq7+0NIk&Ig6YpO%eNI)T{Wts%3%5En&iv9%t%{0XN(sxKcXY8pmTe*pSOM z%e~2;59`v4`Ex9&pFZ0Xx^j#PdTc*ZiS;Ld(} zW;e+urDM2)RwA)HPtsz6(LC8R8n7E@-8fWRhRyjbst#ow-8R?}Gkq>h7UQH(!2i8t zEPD7RbzmLLU`HNTWtn8Jt(vZxTn%~|RJSb2e~OcJ5AY3PMSKdI;xylh2ui^f4953j zwfirthW&77dGAxO9KqAT*o!|=rl2lsu{LyA;MOy$+%&N0wB1ZBc$F-C-RBGIsrJdj zhy19wx)2D5&0{oK(e>LI6*U@Y>lA0L&bWS1d3~`*S!r%Sg$lMNec>xK_|Ix0fz1EL zbtJBu?SLy_&_wAPu*UeHYa1!O8m$;xsV!Moq^1UJX)YyRWLCaI@{6lacL==-`$=Ks zPqz0VfS4@E&+>LQ-z5IGj_@hfAZ5(x$gd`orO(wTjk!Gdxe_#QR;#p-chttk1^9y! OgIU;KDZBi8+&=-q^tC + + + + + + +Cutlass: Member List + + + + + + + + + + +
+
+
+ + +
int const cutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::Rank = kRank
+ + + + + +
+
Cutlass +
+
CUDA Templates for Linear Algebra Subroutines and Solvers
+
+
+ + + + + + + + +
+
+ + +
+ +
+ + +
+
+
+
cutlass::TensorRefArray< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::ConstIterator Member List
+
+
+ +

This is the complete list of members for cutlass::TensorRefArray< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::ConstIterator, including all inherited members.

+ + + + + + + + + + + + +
ConstIterator(TensorArrayRef const &ref, int idx=0)cutlass::TensorRefArray< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::ConstIteratorinline
operator() constcutlass::TensorRefArray< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::ConstIteratorinline
operator+(Index idx)cutlass::TensorRefArray< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::ConstIteratorinline
operator++()cutlass::TensorRefArray< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::ConstIteratorinline
operator++(int)cutlass::TensorRefArray< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::ConstIteratorinline
operator+=(Index idx)cutlass::TensorRefArray< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::ConstIteratorinline
operator-(Index idx)cutlass::TensorRefArray< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::ConstIteratorinline
operator--()cutlass::TensorRefArray< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::ConstIteratorinline
operator--(int)cutlass::TensorRefArray< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::ConstIteratorinline
operator-=(Index idx)cutlass::TensorRefArray< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::ConstIteratorinline
TensorRef typedefcutlass::TensorRefArray< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::ConstIterator
+ +
+Generated on Wed Sep 19 2018 13:58:51 for Cutlass by   +doxygen + 1.8.14 +
+ + diff --git a/docs/classcutlass_1_1TensorRefArray_1_1ConstIterator.html b/docs/classcutlass_1_1TensorRefArray_1_1ConstIterator.html new file mode 100644 index 00000000..aa40085c --- /dev/null +++ b/docs/classcutlass_1_1TensorRefArray_1_1ConstIterator.html @@ -0,0 +1,440 @@ + + + + + + + +Cutlass: cutlass::TensorRefArray< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::ConstIterator Class Reference + + + + + + + + + + +
+
+ + + + + + +
+
Cutlass +
+
CUDA Templates for Linear Algebra Subroutines and Solvers
+
+
+ + + + + + + + +
+
+ + +
+ +
+ + +
+
+
+Public Types | +Public Member Functions | +List of all members
+
+
cutlass::TensorRefArray< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::ConstIterator Class Reference
+
+
+ +

TensorRefIterator over TensorRef objects in TensorRefArray. +

+ +

#include <tensor_ref_collection.h>

+ + + + + +

+Public Types

typedef Base TensorRef
 TensorRef returned by the iterator. More...
 
+ + + + + + + + + + + + + + + + + + + + + + + + + + +

+Public Member Functions

CUTLASS_HOST_DEVICE ConstIterator (TensorArrayRef const &ref, int idx=0)
 Constructs a ConstIterator over the TensorRef objects. More...
 
CUTLASS_HOST_DEVICE TensorRef * operator () const
 Obtains a TensorRef pointed to by this iterator. More...
 
CUTLASS_HOST_DEVICE ConstIterator & operator++ ()
 Advances to next TensorRef. More...
 
CUTLASS_HOST_DEVICE ConstIterator operator++ (int)
 Advances to next TensorRef. More...
 
CUTLASS_HOST_DEVICE ConstIterator operator+ (Index idx)
 
CUTLASS_HOST_DEVICE ConstIterator & operator+= (Index idx)
 
CUTLASS_HOST_DEVICE ConstIterator & operator-- ()
 
CUTLASS_HOST_DEVICE ConstIterator operator-- (int)
 Advances to next TensorRef. More...
 
CUTLASS_HOST_DEVICE ConstIterator & operator-= (Index idx)
 
CUTLASS_HOST_DEVICE ConstIterator operator- (Index idx)
 
+

Member Typedef Documentation

+ +

◆ TensorRef

+ +
+
+
+template<typename Storage_ , int Rank_, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>
+ + + + +
typedef Base cutlass::TensorRefArray< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::ConstIterator::TensorRef
+
+ +
+
+

Constructor & Destructor Documentation

+ +

◆ ConstIterator()

+ +
+
+
+template<typename Storage_ , int Rank_, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>
+ + + + + +
+ + + + + + + + + + + + + + + + + + +
CUTLASS_HOST_DEVICE cutlass::TensorRefArray< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::ConstIterator::ConstIterator (TensorArrayRef const & ref,
int idx = 0 
)
+ 		+inline
+
+ +
+
+

Member Function Documentation

+ +

◆ operator()

+ +
+
+
+template<typename Storage_ , int Rank_, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>
+ + + + + +
+ + + + + + + +
CUTLASS_HOST_DEVICE TensorRef* cutlass::TensorRefArray< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::ConstIterator::operator () const
+ 		+inline
+
+ +
+
+ +

◆ operator+()

+ +
+
+
+template<typename Storage_ , int Rank_, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>
+ + + + + +
+ + + + + + + + +
CUTLASS_HOST_DEVICE ConstIterator cutlass::TensorRefArray< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::ConstIterator::operator+ (Index idx)
+ 		+inline
+
+ +
+
+ +

◆ operator++() [1/2]

+ +
+
+
+template<typename Storage_ , int Rank_, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>
+ + + + + +
+ + + + + + + +
CUTLASS_HOST_DEVICE ConstIterator& cutlass::TensorRefArray< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::ConstIterator::operator++ ()
+ 		+inline
+
+ +
+
+ +

◆ operator++() [2/2]

+ +
+
+
+template<typename Storage_ , int Rank_, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>
+ + + + + +
+ + + + + + + + +
CUTLASS_HOST_DEVICE ConstIterator cutlass::TensorRefArray< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::ConstIterator::operator++ (int )
+ 		+inline
+
+ +
+
+ +

◆ operator+=()

+ +
+
+
+template<typename Storage_ , int Rank_, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>
+ + + + + +
+ + + + + + + + +
CUTLASS_HOST_DEVICE ConstIterator& cutlass::TensorRefArray< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::ConstIterator::operator+= (Index idx)
+ 		+inline
+
+ +
+
+ +

◆ operator-()

+ +
+
+
+template<typename Storage_ , int Rank_, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>
+ + + + + +
+ + + + + + + + +
CUTLASS_HOST_DEVICE ConstIterator cutlass::TensorRefArray< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::ConstIterator::operator- (Index idx)
+ 		+inline
+
+ +
+
+ +

◆ operator--() [1/2]

+ +
+
+
+template<typename Storage_ , int Rank_, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>
+ + + + + +
+ + + + + + + +
CUTLASS_HOST_DEVICE ConstIterator& cutlass::TensorRefArray< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::ConstIterator::operator-- ()
+ 		+inline
+
+ +
+
+ +

◆ operator--() [2/2]

+ +
+
+
+template<typename Storage_ , int Rank_, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>
+ + + + + +
+ + + + + + + + +
CUTLASS_HOST_DEVICE ConstIterator cutlass::TensorRefArray< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::ConstIterator::operator-- (int )
+ 		+inline
+
+ +
+
+ +

◆ operator-=()

+ +
+
+
+template<typename Storage_ , int Rank_, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>
+ + + + + +
+ + + + + + + + +
CUTLASS_HOST_DEVICE ConstIterator& cutlass::TensorRefArray< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::ConstIterator::operator-= (Index idx)
+ 		+inline
+
+ +
+
+
The documentation for this class was generated from the following file: +
+ +
+Generated on Wed Sep 19 2018 13:58:51 for Cutlass by   +doxygen + 1.8.14 +
+ + diff --git a/docs/classcutlass_1_1TensorRefBatchStrided_1_1ConstIterator-members.html b/docs/classcutlass_1_1TensorRefBatchStrided_1_1ConstIterator-members.html new file mode 100644 index 00000000..bb387618 --- /dev/null +++ b/docs/classcutlass_1_1TensorRefBatchStrided_1_1ConstIterator-members.html @@ -0,0 +1,102 @@ + + + + + + + +Cutlass: Member List + + + + + + + + + + +
+
+ + + + + + +
+
Cutlass +
+
CUDA Templates for Linear Algebra Subroutines and Solvers
+
+
+ + + + + + + + +
+
+ + +
+ +
+ + +
+
+
+
cutlass::TensorRefBatchStrided< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::ConstIterator Member List
+
+
+ +

This is the complete list of members for cutlass::TensorRefBatchStrided< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::ConstIterator, including all inherited members.

+ + + + + + + + + + + + + +
ConstIterator(TensorRefBatchStrided const &ref, LongIndex offset=0)cutlass::TensorRefBatchStrided< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::ConstIteratorinline
operator() constcutlass::TensorRefBatchStrided< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::ConstIteratorinline
operator+(Index idx)cutlass::TensorRefBatchStrided< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::ConstIteratorinline
operator++()cutlass::TensorRefBatchStrided< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::ConstIteratorinline
operator++(int)cutlass::TensorRefBatchStrided< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::ConstIteratorinline
operator+=(Index idx)cutlass::TensorRefBatchStrided< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::ConstIteratorinline
operator-(Index idx)cutlass::TensorRefBatchStrided< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::ConstIteratorinline
operator-(ConstIterator const &it)cutlass::TensorRefBatchStrided< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::ConstIteratorinline
operator--()cutlass::TensorRefBatchStrided< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::ConstIteratorinline
operator--(int)cutlass::TensorRefBatchStrided< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::ConstIteratorinline
operator-=(Index idx)cutlass::TensorRefBatchStrided< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::ConstIteratorinline
TensorRef typedefcutlass::TensorRefBatchStrided< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::ConstIterator
+ +
+Generated on Wed Sep 19 2018 13:58:51 for Cutlass by   +doxygen + 1.8.14 +
+ + diff --git a/docs/classcutlass_1_1TensorRefBatchStrided_1_1ConstIterator.html b/docs/classcutlass_1_1TensorRefBatchStrided_1_1ConstIterator.html new file mode 100644 index 00000000..c3dbd9df --- /dev/null +++ b/docs/classcutlass_1_1TensorRefBatchStrided_1_1ConstIterator.html @@ -0,0 +1,476 @@ + + + + + + + +Cutlass: cutlass::TensorRefBatchStrided< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::ConstIterator Class Reference + + + + + + + + + + +
+
+ + + + + + +
+
Cutlass +
+
CUDA Templates for Linear Algebra Subroutines and Solvers
+
+
+ + + + + + + + +
+
+ + +
+ +
+ + +
+
+
+Public Types | +Public Member Functions | +List of all members
+
+
cutlass::TensorRefBatchStrided< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::ConstIterator Class Reference
+
+
+ +

Constant iterator over tensors implied by TensorRefBatchStrided. +

+ +

#include <tensor_ref_collection.h>

+ + + + + +

+Public Types

typedef Base TensorRef
 TensorRef returned by the iterator. More...
 
+ + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + +

+Public Member Functions

CUTLASS_HOST_DEVICE ConstIterator (TensorRefBatchStrided const &ref, LongIndex offset=0)
 Constructs a ConstIterator from a parent TensorRefBatchStrided. More...
 
CUTLASS_HOST_DEVICE TensorRef * operator () const
 Obtains a TensorRef pointed to by the iterator. More...
 
CUTLASS_HOST_DEVICE ConstIterator & operator++ ()
 Advances the iterator to point to the next tensor. More...
 
CUTLASS_HOST_DEVICE ConstIterator operator++ (int)
 Advances the iterator to point to the next tensor. More...
 
CUTLASS_HOST_DEVICE ConstIterator operator+ (Index idx)
 Returns an iterator advanced by (idx) amount. More...
 
CUTLASS_HOST_DEVICE ConstIterator & operator+= (Index idx)
 Advances this iterator by (idx) and returns a reference to self. More...
 
CUTLASS_HOST_DEVICE ConstIterator & operator-- ()
 Moves to the previous tensor. More...
 
CUTLASS_HOST_DEVICE ConstIterator operator-- (int)
 Moves to the previous tensor. More...
 
CUTLASS_HOST_DEVICE ConstIterator operator- (Index idx)
 Returns an iterator moved forward by (idx) amount. More...
 
CUTLASS_HOST_DEVICE ConstIterator & operator-= (Index idx)
 Moves this iterator by (idx) and returns a reference to self. More...
 
CUTLASS_HOST_DEVICE Stride operator- (ConstIterator const &it)
 Returns the difference in offset between two iterators. More...
 
+

Member Typedef Documentation

+ +

◆ TensorRef

+ +
+
+
+template<typename Storage_ , int Rank_, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>
+ + + + +
typedef Base cutlass::TensorRefBatchStrided< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::ConstIterator::TensorRef
+
+ +
+
+

Constructor & Destructor Documentation

+ +

◆ ConstIterator()

+ +
+
+
+template<typename Storage_ , int Rank_, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>
+ + + + + +
+ + + + + + + + + + + + + + + + + + +
CUTLASS_HOST_DEVICE cutlass::TensorRefBatchStrided< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::ConstIterator::ConstIterator (TensorRefBatchStrided const & ref,
LongIndex offset = 0 
)
+ 		+inline
+
+ +
+
+

Member Function Documentation

+ +

◆ operator()

+ +
+
+
+template<typename Storage_ , int Rank_, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>
+ + + + + +
+ + + + + + + +
CUTLASS_HOST_DEVICE TensorRef* cutlass::TensorRefBatchStrided< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::ConstIterator::operator () const
+ 		+inline
+
+ +
+
+ +

◆ operator+()

+ +
+
+
+template<typename Storage_ , int Rank_, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>
+ + + + + +
+ + + + + + + + +
CUTLASS_HOST_DEVICE ConstIterator cutlass::TensorRefBatchStrided< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::ConstIterator::operator+ (Index idx)
+ 		+inline
+
+ +
+
+ +

◆ operator++() [1/2]

+ +
+
+
+template<typename Storage_ , int Rank_, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>
+ + + + + +
+ + + + + + + +
CUTLASS_HOST_DEVICE ConstIterator& cutlass::TensorRefBatchStrided< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::ConstIterator::operator++ ()
+ 		+inline
+
+ +
+
+ +

◆ operator++() [2/2]

+ +
+
+
+template<typename Storage_ , int Rank_, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>
+ + + + + +
+ + + + + + + + +
CUTLASS_HOST_DEVICE ConstIterator cutlass::TensorRefBatchStrided< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::ConstIterator::operator++ (int )
+ 		+inline
+
+ +
+
+ +

◆ operator+=()

+ +
+
+
+template<typename Storage_ , int Rank_, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>
+ + + + + +
+ + + + + + + + +
CUTLASS_HOST_DEVICE ConstIterator& cutlass::TensorRefBatchStrided< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::ConstIterator::operator+= (Index idx)
+ 		+inline
+
+ +
+
+ +

◆ operator-() [1/2]

+ +
+
+
+template<typename Storage_ , int Rank_, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>
+ + + + + +
+ + + + + + + + +
CUTLASS_HOST_DEVICE ConstIterator cutlass::TensorRefBatchStrided< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::ConstIterator::operator- (Index idx)
+ 		+inline
+
+ +
+
+ +

◆ operator-() [2/2]

+ +
+
+
+template<typename Storage_ , int Rank_, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>
+ + + + + +
+ + + + + + + + +
CUTLASS_HOST_DEVICE Stride cutlass::TensorRefBatchStrided< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::ConstIterator::operator- (ConstIterator const & it)
+ 		+inline
+
+ +
+
+ +

◆ operator--() [1/2]

+ +
+
+
+template<typename Storage_ , int Rank_, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>
+ + + + + +
+ + + + + + + +
CUTLASS_HOST_DEVICE ConstIterator& cutlass::TensorRefBatchStrided< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::ConstIterator::operator-- ()
+ 		+inline
+
+ +
+
+ +

◆ operator--() [2/2]

+ +
+
+
+template<typename Storage_ , int Rank_, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>
+ + + + + +
+ + + + + + + + +
CUTLASS_HOST_DEVICE ConstIterator cutlass::TensorRefBatchStrided< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::ConstIterator::operator-- (int )
+ 		+inline
+
+ +
+
+ +

◆ operator-=()

+ +
+
+
+template<typename Storage_ , int Rank_, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>
+ + + + + +
+ + + + + + + + +
CUTLASS_HOST_DEVICE ConstIterator& cutlass::TensorRefBatchStrided< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::ConstIterator::operator-= (Index idx)
+ 		+inline
+
+ +
+
+
The documentation for this class was generated from the following file: +
+ +
+Generated on Wed Sep 19 2018 13:58:51 for Cutlass by   +doxygen + 1.8.14 +
+ + diff --git a/docs/classcutlass_1_1TensorRef_3_01Storage___00_01Rank___00_01MapFunc___00_011_00_01Index___00_01LongIndex___01_4-members.html b/docs/classcutlass_1_1TensorRef_3_01Storage___00_01Rank___00_01MapFunc___00_011_00_01Index___00_01LongIndex___01_4-members.html new file mode 100644 index 00000000..8af74ab9 --- /dev/null +++ b/docs/classcutlass_1_1TensorRef_3_01Storage___00_01Rank___00_01MapFunc___00_011_00_01Index___00_01LongIndex___01_4-members.html @@ -0,0 +1,124 @@ + + + + + + + +Cutlass: Member List + + + + + + + + + + +
+
+ + + + + + +
+
Cutlass +
+
CUDA Templates for Linear Algebra Subroutines and Solvers
+
+
+ + + + + + + + +
+
+ + +
+ +
+ + +
+
+
+
cutlass::TensorRef< Storage_, Rank_, MapFunc_, 1, Index_, LongIndex_ > Member List
+
+
+ +

This is the complete list of members for cutlass::TensorRef< Storage_, Rank_, MapFunc_, 1, Index_, LongIndex_ >, including all inherited members.

+ + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + +
add_pointer_offset(LongIndex delta)cutlass::TensorRef< Storage_, Rank_, MapFunc_, 1, Index_, LongIndex_ >inline
at(TensorCoord const &coord) constcutlass::TensorRef< Storage_, Rank_, MapFunc_, 1, Index_, LongIndex_ >inline
at(LongIndex idx) constcutlass::TensorRef< Storage_, Rank_, MapFunc_, 1, Index_, LongIndex_ >inline
const_ref() constcutlass::TensorRef< Storage_, Rank_, MapFunc_, 1, Index_, LongIndex_ >inline
ConstTensorRef typedefcutlass::TensorRef< Storage_, Rank_, MapFunc_, 1, Index_, LongIndex_ >
Coord_t typedefcutlass::TensorRef< Storage_, Rank_, MapFunc_, 1, Index_, LongIndex_ >
data() constcutlass::TensorRef< Storage_, Rank_, MapFunc_, 1, Index_, LongIndex_ >inline
good() constcutlass::TensorRef< Storage_, Rank_, MapFunc_, 1, Index_, LongIndex_ >inline
Index typedefcutlass::TensorRef< Storage_, Rank_, MapFunc_, 1, Index_, LongIndex_ >
kRankcutlass::TensorRef< Storage_, Rank_, MapFunc_, 1, Index_, LongIndex_ >static
kStorageRankcutlass::TensorRef< Storage_, Rank_, MapFunc_, 1, Index_, LongIndex_ >static
leading_dim(int idx=0) constcutlass::TensorRef< Storage_, Rank_, MapFunc_, 1, Index_, LongIndex_ >inline
LongIndex typedefcutlass::TensorRef< Storage_, Rank_, MapFunc_, 1, Index_, LongIndex_ >
map(TensorCoord const &coord) constcutlass::TensorRef< Storage_, Rank_, MapFunc_, 1, Index_, LongIndex_ >inline
MapFunc typedefcutlass::TensorRef< Storage_, Rank_, MapFunc_, 1, Index_, LongIndex_ >
offset(TensorCoord const &coord) constcutlass::TensorRef< Storage_, Rank_, MapFunc_, 1, Index_, LongIndex_ >inline
operator+(TensorCoord const &b) constcutlass::TensorRef< Storage_, Rank_, MapFunc_, 1, Index_, LongIndex_ >inline
operator+=(TensorCoord const &b)cutlass::TensorRef< Storage_, Rank_, MapFunc_, 1, Index_, LongIndex_ >inline
operator-(TensorCoord const &b) constcutlass::TensorRef< Storage_, Rank_, MapFunc_, 1, Index_, LongIndex_ >inline
operator-=(TensorCoord const &b)cutlass::TensorRef< Storage_, Rank_, MapFunc_, 1, Index_, LongIndex_ >inline
operator[](TensorCoord const &coord) constcutlass::TensorRef< Storage_, Rank_, MapFunc_, 1, Index_, LongIndex_ >inline
operator[](LongIndex idx) constcutlass::TensorRef< Storage_, Rank_, MapFunc_, 1, Index_, LongIndex_ >inline
Rankcutlass::TensorRef< Storage_, Rank_, MapFunc_, 1, Index_, LongIndex_ >static
reset(Storage *ptr=nullptr)cutlass::TensorRef< Storage_, Rank_, MapFunc_, 1, Index_, LongIndex_ >inline
reset(Storage *ptr, StorageCoord const &stride)cutlass::TensorRef< Storage_, Rank_, MapFunc_, 1, Index_, LongIndex_ >inline
Storage typedefcutlass::TensorRef< Storage_, Rank_, MapFunc_, 1, Index_, LongIndex_ >
StorageCoord typedefcutlass::TensorRef< Storage_, Rank_, MapFunc_, 1, Index_, LongIndex_ >
stride() constcutlass::TensorRef< Storage_, Rank_, MapFunc_, 1, Index_, LongIndex_ >inline
stride(int dim) constcutlass::TensorRef< Storage_, Rank_, MapFunc_, 1, Index_, LongIndex_ >inline
TensorCoord typedefcutlass::TensorRef< Storage_, Rank_, MapFunc_, 1, Index_, LongIndex_ >
TensorRef(Storage *ptr=nullptr)cutlass::TensorRef< Storage_, Rank_, MapFunc_, 1, Index_, LongIndex_ >inline
TensorRef(Storage *ptr, StrideVector const &stride)cutlass::TensorRef< Storage_, Rank_, MapFunc_, 1, Index_, LongIndex_ >inline
TensorRef(Storage *ptr, StorageCoord const &stride)cutlass::TensorRef< Storage_, Rank_, MapFunc_, 1, Index_, LongIndex_ >inline
TensorRef(TensorRef< typename platform::remove_const< Storage >::type, kRank, MapFunc, kStorageRank, Index, LongIndex > const &ref)cutlass::TensorRef< Storage_, Rank_, MapFunc_, 1, Index_, LongIndex_ >inline
+ +
+Generated on Wed Sep 19 2018 13:58:51 for Cutlass by   +doxygen + 1.8.14 +
+ + diff --git a/docs/classcutlass_1_1TensorRef_3_01Storage___00_01Rank___00_01MapFunc___00_011_00_01Index___00_01LongIndex___01_4.html b/docs/classcutlass_1_1TensorRef_3_01Storage___00_01Rank___00_01MapFunc___00_011_00_01Index___00_01LongIndex___01_4.html new file mode 100644 index 00000000..2dfd10c9 --- /dev/null +++ b/docs/classcutlass_1_1TensorRef_3_01Storage___00_01Rank___00_01MapFunc___00_011_00_01Index___00_01LongIndex___01_4.html @@ -0,0 +1,1092 @@ + + + + + + + +Cutlass: cutlass::TensorRef< Storage_, Rank_, MapFunc_, 1, Index_, LongIndex_ > Class Template Reference + + + + + + + + + + +
+
+ + + + + + +
+
Cutlass +
+
CUDA Templates for Linear Algebra Subroutines and Solvers
+
+
+ + + + + + + + +
+
+ + +
+ +
+ + +
+
+
+Classes | +Public Types | +Public Member Functions | +Static Public Attributes | +List of all members
+
+
cutlass::TensorRef< Storage_, Rank_, MapFunc_, 1, Index_, LongIndex_ > Class Template Reference
+
+
+ +

Specialization for rank=1 case with no internal StrideVector. +

+ +

#include <tensor_ref.h>

+ + + + +

+Classes

struct  StrideVector
 
+ + + + + + + + + + + + + + + + + + + + + + + + + +

+Public Types

typedef Storage_ Storage
 Data type of individual access. More...
 
typedef MapFunc_ MapFunc
 Mapping function from logical coordinate to internal n-D array. More...
 
typedef Index_ Index
 Index type. More...
 
typedef LongIndex_ LongIndex
 Typically, strides in memory can be very large. More...
 
typedef Coord< kRank > TensorCoord
 Coordinate in logical tensor space. More...
 
typedef Coord< kStorageRank > StorageCoord
 Coordinate in storage n-D array. More...
 
typedef TensorRef< typename platform::remove_const< Storage >::type const, Rank_, MapFunc_, kStorageRank, Index_, LongIndex_ > ConstTensorRef
 Tensor reference to of constant value. More...
 
typedef TensorCoord Coord_t
 Coordinate in logical tensor space. More...
 
+ + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + +

+Public Member Functions

CUTLASS_HOST_DEVICE TensorRef (Storage *ptr=nullptr)
 Helper for 1-D memory. All higher ranks are projected onto the fastest changing rank. More...
 
CUTLASS_HOST_DEVICE TensorRef (Storage *ptr, StrideVector const &stride)
 Constructs from a single pointer and stride vector. More...
 
CUTLASS_HOST_DEVICE TensorRef (Storage *ptr, StorageCoord const &stride)
 
CUTLASS_HOST_DEVICE TensorRef (TensorRef< typename platform::remove_const< Storage >::type, kRank, MapFunc, kStorageRank, Index, LongIndex > const &ref)
 Enables conversion from TensorRef of non-const type. More...
 
CUTLASS_HOST_DEVICE ConstTensorRef const_ref () const
 Returns a reference to constant-valued tensor. More...
 
CUTLASS_HOST_DEVICE void reset (Storage *ptr=nullptr)
 Updates only the pointer. More...
 
CUTLASS_HOST_DEVICE void reset (Storage *ptr, StorageCoord const &stride)
 Updates the pointer, stride, and location within a TensorRef. More...
 
CUTLASS_HOST_DEVICE bool good () const
 Returns true if the TensorRef may be safely accessed. More...
 
CUTLASS_HOST_DEVICE Storage * data () const
 Returns the pointer to referenced data. More...
 
CUTLASS_HOST_DEVICE StorageCoord stride () const
 Returns the stride of the tensor. More...
 
CUTLASS_HOST_DEVICE Index stride (int dim) const
 Returns the stride of the tensor in the given dimension. More...
 
CUTLASS_HOST_DEVICE Index leading_dim (int idx=0) const
 Returns the maximum stride element as the 'leading dimension'. More...
 
CUTLASS_HOST_DEVICE StorageCoord map (TensorCoord const &coord) const
 Maps a logical coordinate to an n-D array in memory. More...
 
CUTLASS_HOST_DEVICE LongIndex offset (TensorCoord const &coord) const
 Computes the offset of an index from the origin of the tensor. More...
 
CUTLASS_HOST_DEVICE Storage & at (TensorCoord const &coord) const
 Returns a reference to the element at a given Coord. More...
 
CUTLASS_HOST_DEVICE Storage & at (LongIndex idx) const
 Returns a reference to the element at a given linear index. More...
 
CUTLASS_HOST_DEVICE Storage & operator[] (TensorCoord const &coord) const
 Returns a reference to the element at a given Coord. More...
 
CUTLASS_HOST_DEVICE Storage & operator[] (LongIndex idx) const
 Returns a reference to the element at a given linear index. More...
 
CUTLASS_HOST_DEVICE TensorRef & add_pointer_offset (LongIndex delta)
 Adds an offset to each pointer. More...
 
CUTLASS_HOST_DEVICE TensorRef operator+ (TensorCoord const &b) const
 Returns a TensorRef offset by a given amount. More...
 
CUTLASS_HOST_DEVICE TensorRef & operator+= (TensorCoord const &b)
 Returns a TensorRef offset by a given amount. More...
 
CUTLASS_HOST_DEVICE TensorRef operator- (TensorCoord const &b) const
 Returns a TensorRef offset by a given amount. More...
 
CUTLASS_HOST_DEVICE TensorRef & operator-= (TensorCoord const &b)
 Returns a TensorRef offset by a given amount. More...
 
+ + + + + + + + + + +

+Static Public Attributes

static int const kRank = Rank_
 Logical rank of tensor index space. More...
 
static int const kStorageRank = 1
 Rank of internal storage. More...
 
static int const Rank = kRank
 Logical rank of tensor index space. More...
 
+

Member Typedef Documentation

+ +

◆ ConstTensorRef

+ +
+
+
+template<typename Storage_ , int Rank_, typename MapFunc_ , typename Index_ , typename LongIndex_ >
+ + + + +
typedef TensorRef< typename platform::remove_const<Storage>::type const, Rank_, MapFunc_, kStorageRank, Index_, LongIndex_> cutlass::TensorRef< Storage_, Rank_, MapFunc_, 1, Index_, LongIndex_ >::ConstTensorRef
+
+ +
+
+ +

◆ Coord_t

+ +
+
+
+template<typename Storage_ , int Rank_, typename MapFunc_ , typename Index_ , typename LongIndex_ >
+ + + + +
typedef TensorCoord cutlass::TensorRef< Storage_, Rank_, MapFunc_, 1, Index_, LongIndex_ >::Coord_t
+
+ +
+
+ +

◆ Index

+ +
+
+
+template<typename Storage_ , int Rank_, typename MapFunc_ , typename Index_ , typename LongIndex_ >
+ + + + +
typedef Index_ cutlass::TensorRef< Storage_, Rank_, MapFunc_, 1, Index_, LongIndex_ >::Index
+
+ +
+
+ +

◆ LongIndex

+ +
+
+
+template<typename Storage_ , int Rank_, typename MapFunc_ , typename Index_ , typename LongIndex_ >
+ + + + +
typedef LongIndex_ cutlass::TensorRef< Storage_, Rank_, MapFunc_, 1, Index_, LongIndex_ >::LongIndex
+
+ +
+
+ +

◆ MapFunc

+ +
+
+
+template<typename Storage_ , int Rank_, typename MapFunc_ , typename Index_ , typename LongIndex_ >
+ + + + +
typedef MapFunc_ cutlass::TensorRef< Storage_, Rank_, MapFunc_, 1, Index_, LongIndex_ >::MapFunc
+
+ +
+
+ +

◆ Storage

+ +
+
+
+template<typename Storage_ , int Rank_, typename MapFunc_ , typename Index_ , typename LongIndex_ >
+ + + + +
typedef Storage_ cutlass::TensorRef< Storage_, Rank_, MapFunc_, 1, Index_, LongIndex_ >::Storage
+
+ +
+
+ +

◆ StorageCoord

+ +
+
+
+template<typename Storage_ , int Rank_, typename MapFunc_ , typename Index_ , typename LongIndex_ >
+ + + + +
typedef Coord<kStorageRank> cutlass::TensorRef< Storage_, Rank_, MapFunc_, 1, Index_, LongIndex_ >::StorageCoord
+
+ +
+
+ +

◆ TensorCoord

+ +
+
+
+template<typename Storage_ , int Rank_, typename MapFunc_ , typename Index_ , typename LongIndex_ >
+ + + + +
typedef Coord<kRank> cutlass::TensorRef< Storage_, Rank_, MapFunc_, 1, Index_, LongIndex_ >::TensorCoord
+
+ +
+
+

Constructor & Destructor Documentation

+ +

◆ TensorRef() [1/4]

+ +
+
+
+template<typename Storage_ , int Rank_, typename MapFunc_ , typename Index_ , typename LongIndex_ >
+ + + + + +
+ + + + + + + + +
CUTLASS_HOST_DEVICE cutlass::TensorRef< Storage_, Rank_, MapFunc_, 1, Index_, LongIndex_ >::TensorRef (Storage * ptr = nullptr)
+ 		+inline
+
+ +
+
+ +

◆ TensorRef() [2/4]

+ +
+
+
+template<typename Storage_ , int Rank_, typename MapFunc_ , typename Index_ , typename LongIndex_ >
+ + + + + +
+ + + + + + + + + + + + + + + + + + +
CUTLASS_HOST_DEVICE cutlass::TensorRef< Storage_, Rank_, MapFunc_, 1, Index_, LongIndex_ >::TensorRef (Storage * ptr,
StrideVector const & stride 
)
+ 		+inline
+
+ +
+
+ +

◆ TensorRef() [3/4]

+ +
+
+
+template<typename Storage_ , int Rank_, typename MapFunc_ , typename Index_ , typename LongIndex_ >
+ + + + + +
+ + + + + + + + + + + + + + + + + + +
CUTLASS_HOST_DEVICE cutlass::TensorRef< Storage_, Rank_, MapFunc_, 1, Index_, LongIndex_ >::TensorRef (Storage * ptr,
StorageCoord const & stride 
)
+ 		+inline
+
+

Constructs from a pointer and a stride vector of size kRank. If fastest changing stride is not 1, construction fails and subsequent calls to good() will return false.

+ +
+
+ +

◆ TensorRef() [4/4]

+ +
+
+
+template<typename Storage_ , int Rank_, typename MapFunc_ , typename Index_ , typename LongIndex_ >
+ + + + + +
+ + + + + + + + +
CUTLASS_HOST_DEVICE cutlass::TensorRef< Storage_, Rank_, MapFunc_, 1, Index_, LongIndex_ >::TensorRef (TensorRef< typename platform::remove_const< Storage >::type, kRank, MapFunc, kStorageRank, Index, LongIndex > const & ref)
+ 		+inline
+
+ +
+
+

Member Function Documentation

+ +

◆ add_pointer_offset()

+ +
+
+
+template<typename Storage_ , int Rank_, typename MapFunc_ , typename Index_ , typename LongIndex_ >
+ + + + + +
+ + + + + + + + +
CUTLASS_HOST_DEVICE TensorRef& cutlass::TensorRef< Storage_, Rank_, MapFunc_, 1, Index_, LongIndex_ >::add_pointer_offset (LongIndex delta)
+ 		+inline
+
+ +
+
+ +

◆ at() [1/2]

+ +
+
+
+template<typename Storage_ , int Rank_, typename MapFunc_ , typename Index_ , typename LongIndex_ >
+ + + + + +
+ + + + + + + + +
CUTLASS_HOST_DEVICE Storage& cutlass::TensorRef< Storage_, Rank_, MapFunc_, 1, Index_, LongIndex_ >::at (TensorCoord const & coord) const
+ 		+inline
+
+ +
+
+ +

◆ at() [2/2]

+ +
+
+
+template<typename Storage_ , int Rank_, typename MapFunc_ , typename Index_ , typename LongIndex_ >
+ + + + + +
+ + + + + + + + +
CUTLASS_HOST_DEVICE Storage& cutlass::TensorRef< Storage_, Rank_, MapFunc_, 1, Index_, LongIndex_ >::at (LongIndex idx) const
+ 		+inline
+
+ +
+
+ +

◆ const_ref()

+ +
+
+
+template<typename Storage_ , int Rank_, typename MapFunc_ , typename Index_ , typename LongIndex_ >
+ + + + + +
+ + + + + + + +
CUTLASS_HOST_DEVICE ConstTensorRef cutlass::TensorRef< Storage_, Rank_, MapFunc_, 1, Index_, LongIndex_ >::const_ref () const
+ 		+inline
+
+ +
+
+ +

◆ data()

+ +
+
+
+template<typename Storage_ , int Rank_, typename MapFunc_ , typename Index_ , typename LongIndex_ >
+ + + + + +
+ + + + + + + +
CUTLASS_HOST_DEVICE Storage* cutlass::TensorRef< Storage_, Rank_, MapFunc_, 1, Index_, LongIndex_ >::data () const
+ 		+inline
+
+ +
+
+ +

◆ good()

+ +
+
+
+template<typename Storage_ , int Rank_, typename MapFunc_ , typename Index_ , typename LongIndex_ >
+ + + + + +
+ + + + + + + +
CUTLASS_HOST_DEVICE bool cutlass::TensorRef< Storage_, Rank_, MapFunc_, 1, Index_, LongIndex_ >::good () const
+ 		+inline
+
+ +
+
+ +

◆ leading_dim()

+ +
+
+
+template<typename Storage_ , int Rank_, typename MapFunc_ , typename Index_ , typename LongIndex_ >
+ + + + + +
+ + + + + + + + +
CUTLASS_HOST_DEVICE Index cutlass::TensorRef< Storage_, Rank_, MapFunc_, 1, Index_, LongIndex_ >::leading_dim (int idx = 0) const
+ 		+inline
+
+ +
+
+ +

◆ map()

+ +
+
+
+template<typename Storage_ , int Rank_, typename MapFunc_ , typename Index_ , typename LongIndex_ >
+ + + + + +
+ + + + + + + + +
CUTLASS_HOST_DEVICE StorageCoord cutlass::TensorRef< Storage_, Rank_, MapFunc_, 1, Index_, LongIndex_ >::map (TensorCoord const & coord) const
+ 		+inline
+
+ +
+
+ +

◆ offset()

+ +
+
+
+template<typename Storage_ , int Rank_, typename MapFunc_ , typename Index_ , typename LongIndex_ >
+ + + + + +
+ + + + + + + + +
CUTLASS_HOST_DEVICE LongIndex cutlass::TensorRef< Storage_, Rank_, MapFunc_, 1, Index_, LongIndex_ >::offset (TensorCoord const & coord) const
+ 		+inline
+
+ +
+
+ +

◆ operator+()

+ +
+
+
+template<typename Storage_ , int Rank_, typename MapFunc_ , typename Index_ , typename LongIndex_ >
+ + + + + +
+ + + + + + + + +
CUTLASS_HOST_DEVICE TensorRef cutlass::TensorRef< Storage_, Rank_, MapFunc_, 1, Index_, LongIndex_ >::operator+ (TensorCoord const & b) const
+ 		+inline
+
+ +
+
+ +

◆ operator+=()

+ +
+
+
+template<typename Storage_ , int Rank_, typename MapFunc_ , typename Index_ , typename LongIndex_ >
+ + + + + +
+ + + + + + + + +
CUTLASS_HOST_DEVICE TensorRef& cutlass::TensorRef< Storage_, Rank_, MapFunc_, 1, Index_, LongIndex_ >::operator+= (TensorCoord const & b)
+ 		+inline
+
+ +
+
+ +

◆ operator-()

+ +
+
+
+template<typename Storage_ , int Rank_, typename MapFunc_ , typename Index_ , typename LongIndex_ >
+ + + + + +
+ + + + + + + + +
CUTLASS_HOST_DEVICE TensorRef cutlass::TensorRef< Storage_, Rank_, MapFunc_, 1, Index_, LongIndex_ >::operator- (TensorCoord const & b) const
+ 		+inline
+
+ +
+
+ +

◆ operator-=()

+ +
+
+
+template<typename Storage_ , int Rank_, typename MapFunc_ , typename Index_ , typename LongIndex_ >
+ + + + + +
+ + + + + + + + +
CUTLASS_HOST_DEVICE TensorRef& cutlass::TensorRef< Storage_, Rank_, MapFunc_, 1, Index_, LongIndex_ >::operator-= (TensorCoord const & b)
+ 		+inline
+
+ +
+
+ +

◆ operator[]() [1/2]

+ +
+
+
+template<typename Storage_ , int Rank_, typename MapFunc_ , typename Index_ , typename LongIndex_ >
+ + + + + +
+ + + + + + + + +
CUTLASS_HOST_DEVICE Storage& cutlass::TensorRef< Storage_, Rank_, MapFunc_, 1, Index_, LongIndex_ >::operator[] (TensorCoord const & coord) const
+ 		+inline
+
+ +
+
+ +

◆ operator[]() [2/2]

+ +
+
+
+template<typename Storage_ , int Rank_, typename MapFunc_ , typename Index_ , typename LongIndex_ >
+ + + + + +
+ + + + + + + + +
CUTLASS_HOST_DEVICE Storage& cutlass::TensorRef< Storage_, Rank_, MapFunc_, 1, Index_, LongIndex_ >::operator[] (LongIndex idx) const
+ 		+inline
+
+ +
+
+ +

◆ reset() [1/2]

+ +
+
+
+template<typename Storage_ , int Rank_, typename MapFunc_ , typename Index_ , typename LongIndex_ >
+ + + + + +
+ + + + + + + + +
CUTLASS_HOST_DEVICE void cutlass::TensorRef< Storage_, Rank_, MapFunc_, 1, Index_, LongIndex_ >::reset (Storage * ptr = nullptr)
+ 		+inline
+
+ +
+
+ +

◆ reset() [2/2]

+ +
+
+
+template<typename Storage_ , int Rank_, typename MapFunc_ , typename Index_ , typename LongIndex_ >
+ + + + + +
+ + + + + + + + + + + + + + + + + + +
CUTLASS_HOST_DEVICE void cutlass::TensorRef< Storage_, Rank_, MapFunc_, 1, Index_, LongIndex_ >::reset (Storage * ptr,
StorageCoord const & stride 
)
+ 		+inline
+
+ +
+
+ +

◆ stride() [1/2]

+ +
+
+
+template<typename Storage_ , int Rank_, typename MapFunc_ , typename Index_ , typename LongIndex_ >
+ + + + + +
+ + + + + + + +
CUTLASS_HOST_DEVICE StorageCoord cutlass::TensorRef< Storage_, Rank_, MapFunc_, 1, Index_, LongIndex_ >::stride () const
+ 		+inline
+
+ +
+
+ +

◆ stride() [2/2]

+ +
+
+
+template<typename Storage_ , int Rank_, typename MapFunc_ , typename Index_ , typename LongIndex_ >
+ + + + + +
+ + + + + + + + +
CUTLASS_HOST_DEVICE Index cutlass::TensorRef< Storage_, Rank_, MapFunc_, 1, Index_, LongIndex_ >::stride (int dim) const
+ 		+inline
+
+ +
+
+

Member Data Documentation

+ +

◆ kRank

+ +
+
+
+template<typename Storage_ , int Rank_, typename MapFunc_ , typename Index_ , typename LongIndex_ >
+ + + + + +
+ + + + +
int const cutlass::TensorRef< Storage_, Rank_, MapFunc_, 1, Index_, LongIndex_ >::kRank = Rank_
+ 		+static
+
+ +
+
+ +

◆ kStorageRank

+ +
+
+
+template<typename Storage_ , int Rank_, typename MapFunc_ , typename Index_ , typename LongIndex_ >
+ + + + + +
+ + + + +
int const cutlass::TensorRef< Storage_, Rank_, MapFunc_, 1, Index_, LongIndex_ >::kStorageRank = 1
+ 		+static
+
+ +
+
+ +

◆ Rank

+ +
+
+
+template<typename Storage_ , int Rank_, typename MapFunc_ , typename Index_ , typename LongIndex_ >
+ + + + + +
+ + + + +
int const cutlass::TensorRef< Storage_, Rank_, MapFunc_, 1, Index_, LongIndex_ >::Rank = kRank
+ 		+static
+
+ +
+
+
The documentation for this class was generated from the following file: +
+ +
+Generated on Wed Sep 19 2018 13:58:51 for Cutlass by   +doxygen + 1.8.14 +
+ + diff --git a/docs/classcutlass_1_1TensorView-members.html b/docs/classcutlass_1_1TensorView-members.html index e9401f9c..9f5c3253 100644 --- a/docs/classcutlass_1_1TensorView-members.html +++ b/docs/classcutlass_1_1TensorView-members.html @@ -73,51 +73,70 @@ $(function() {
-
cutlass::TensorView< T > Member List
+
cutlass::TensorView< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ > Member List
-

This is the complete list of members for cutlass::TensorView< T >, including all inherited members.

+

This is the complete list of members for cutlass::TensorView< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >, including all inherited members.

- - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + +
advance(Coord< Rank > const &b)cutlass::TensorRef< T, 4 >inline
at(Coord_t const &coord) constcutlass::TensorView< T >inline
at(Offset_t idx) constcutlass::TensorView< T >inline
Base typedefcutlass::TensorView< T >
const_ref()cutlass::TensorView< T >inline
ConstTensorRef_t typedefcutlass::TensorView< T >
contains(Coord_t const &coord) constcutlass::TensorView< T >inline
convert()cutlass::TensorRef< T, 4 >inline
Coord_t typedefcutlass::TensorView< T >
data() constcutlass::TensorView< T >inline
good() constcutlass::TensorView< T >inline
leading_dim() constcutlass::TensorRef< T, 4 >inline
offset(Coord_t const &coord) constcutlass::TensorView< T >inline
Offset_t typedefcutlass::TensorView< T >
operator+(Coord< Rank > const &b) constcutlass::TensorRef< T, 4 >inline
operator-(Coord< Rank > const &b) constcutlass::TensorRef< T, 4 >inline
operator=(TensorView const &_tensor)cutlass::TensorView< T >inline
operator[](Coord< Rank > const &coord) constcutlass::TensorView< T >inline
TensorRef< T, 4 >::operator[](int idx) constcutlass::TensorRef< T, 4 >inline
Rankcutlass::TensorView< T >static
ref()cutlass::TensorView< T >inline
ref() constcutlass::TensorView< T >inline
reset(TensorRef_t const &_ref=TensorRef_t(0), Coord_t const &_size=Coord_t())cutlass::TensorView< T >inline
TensorRef< T, 4 >::reset(Storage *ptr=nullptr, Coord< Rank > stride=Coord< Rank >(0))cutlass::TensorRef< T, 4 >inline
size() constcutlass::TensorView< T >inline
size(int dim) constcutlass::TensorView< T >inline
Storage typedefcutlass::TensorRef< T, 4 >
stride() constcutlass::TensorView< T >inline
stride(int dim) constcutlass::TensorView< T >inline
subview(Coord_t const &location, Coord_t size) constcutlass::TensorView< T >inline
TensorRef()cutlass::TensorRef< T, 4 >inline
TensorRef(Storage *ptr, Coord< Rank > stride)cutlass::TensorRef< T, 4 >inline
TensorRef_t typedefcutlass::TensorView< T >
TensorView()cutlass::TensorView< T >inline
TensorView(TensorRef_t const &_ref, Coord_t const &_size)cutlass::TensorView< T >inline
add_pointer_offset(LongIndex delta)cutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >inline
at(TensorCoord const &coord) constcutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >inline
at(LongIndex idx) constcutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >inline
Base typedefcutlass::TensorView< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >
capacity() constcutlass::TensorView< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >inline
const_ref() constcutlass::TensorView< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >inline
ConstTensorRef typedefcutlass::TensorView< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >
ConstTensorRef_t typedefcutlass::TensorView< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >
ConstTensorView typedefcutlass::TensorView< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >
contains(TensorCoord const &coord) constcutlass::TensorView< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >inline
Coord_t typedefcutlass::TensorView< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >
data() constcutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >inline
good() constcutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >inline
Index typedefcutlass::TensorView< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >
kRankcutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >static
kStorageRankcutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >static
leading_dim(int idx=0) constcutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >inline
LongIndex typedefcutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >
map(TensorCoord const &coord) constcutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >inline
MapFunc typedefcutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >
offset(TensorCoord const &coord) constcutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >inline
Offset_t typedefcutlass::TensorView< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >
operator+(TensorCoord const &b) constcutlass::TensorView< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >inline
operator+=(TensorCoord const &b)cutlass::TensorView< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >inline
operator-(TensorCoord const &b) constcutlass::TensorView< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >inline
operator-=(TensorCoord const &b)cutlass::TensorView< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >inline
operator=(TensorView const &_tensor)cutlass::TensorView< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >inline
operator[](TensorCoord const &coord) constcutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >inline
operator[](LongIndex idx) constcutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >inline
Rankcutlass::TensorView< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >static
ref() constcutlass::TensorView< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >inline
reset(Base const &_ref=Base(), TensorCoord const &_size=TensorCoord())cutlass::TensorView< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >inline
cutlass::TensorRef::reset(Storage *ptr=nullptr)cutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >inline
cutlass::TensorRef::reset(Storage *ptr, StorageCoord const &stride)cutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >inline
size() constcutlass::TensorView< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >inline
size(int dim) constcutlass::TensorView< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >inline
Storage typedefcutlass::TensorView< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >
StorageCoord typedefcutlass::TensorView< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >
stride() constcutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >inline
stride(int dim) constcutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >inline
StrideVector typedefcutlass::TensorView< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >
subview(TensorCoord const &location, TensorCoord size) constcutlass::TensorView< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >inline
TensorCoord typedefcutlass::TensorView< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >
TensorRef typedefcutlass::TensorView< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >
cutlass::TensorRef::TensorRef(Storage *ptr=nullptr)cutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >inline
cutlass::TensorRef::TensorRef(Storage *ptr, Index ldm)cutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >inline
cutlass::TensorRef::TensorRef(Storage *ptr, StrideVector const &stride)cutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >inline
cutlass::TensorRef::TensorRef(Storage *ptr, StorageCoord const &stride)cutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >inline
cutlass::TensorRef::TensorRef(TensorRef< typename platform::remove_const< Storage >::type, kRank, MapFunc, kStorageRank, Index, LongIndex > const &ref)cutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >inline
TensorRef_t typedefcutlass::TensorView< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >
TensorView()cutlass::TensorView< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >inline
TensorView(Base const &_ref, TensorCoord const &_size)cutlass::TensorView< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >inline
TensorView(Storage *ptr, StrideVector const &stride, TensorCoord const &size)cutlass::TensorView< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >inline
TensorView(Storage *ptr, StorageCoord const &stride, TensorCoord const &size)cutlass::TensorView< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >inline
-Generated on Thu May 3 2018 16:36:13 for Cutlass by   +Generated on Wed Sep 19 2018 13:58:51 for Cutlass by   doxygen 1.8.14
diff --git a/docs/classcutlass_1_1TensorView.html b/docs/classcutlass_1_1TensorView.html index 7dba2322..276d1077 100644 --- a/docs/classcutlass_1_1TensorView.html +++ b/docs/classcutlass_1_1TensorView.html @@ -5,7 +5,7 @@ -Cutlass: cutlass::TensorView< T > Class Template Reference +Cutlass: cutlass::TensorView< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ > Class Template Reference @@ -78,242 +78,438 @@ $(function() { Static Public Attributes | List of all members
-
cutlass::TensorView< T > Class Template Reference
+
cutlass::TensorView< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ > Class Template Reference
-

Host-side reference implementation of tensor operations. +

Defines a view into a logical tensor.

#include <tensor_view.h>

-Inheritance diagram for cutlass::TensorView< T >:
+Inheritance diagram for cutlass::TensorView< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >:
- - -cutlass::TensorRef< T, 4 > + + +cutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >
- - - - - - - - - - - - - - - - - - - + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + +

Public Types

typedef TensorRef< T, 4 > Base
 Reference and stride. More...
 
typedef Base TensorRef_t
 Reference and stride. More...
 
typedef TensorRef< T const, 4 > ConstTensorRef_t
 Reference to constant type. More...
 
typedef int Offset_t
 Type used to compute the offset of an element to the base of a tensor. More...
 
typedef Coord< Rank > Coord_t
 Coordinate into tensor. More...
 
- Public Types inherited from cutlass::TensorRef< T, 4 >
typedef T Storage
 Data type of individual access. More...
 
typedef TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ > Base
 Base tensor reference. More...
 
typedef TensorRef< typename platform::remove_const< Storage_ >::type const, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ > ConstTensorRef
 Tensor reference to of constant value. More...
 
typedef Base TensorRef
 Base tensor reference. More...
 
typedef Base::Storage Storage
 Storage type. More...
 
typedef Base::Index Index
 Index type. More...
 
typedef TensorRef::TensorCoord TensorCoord
 Coordinate in logical tensor space. More...
 
typedef TensorRef::StorageCoord StorageCoord
 Coordinate in storage n-D array. More...
 
typedef TensorRef::StrideVector StrideVector
 
typedef TensorView< typename platform::remove_const< Storage >::type const, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ > ConstTensorView
 TensorView of constant value. More...
 
typedef TensorCoord Coord_t
 Coordinate in logical tensor space. More...
 
typedef Base::LongIndex Offset_t
 Type used to compute the offset of an element to the base of a tensor. More...
 
typedef TensorRef TensorRef_t
 Base class. More...
 
typedef TensorRef::ConstTensorRef ConstTensorRef_t
 TensorRef to const-valued type. More...
 
- Public Types inherited from cutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >
typedef Storage_ Storage
 Data type of individual access. More...
 
typedef MapFunc_ MapFunc
 Mapping function from logical coordinate to internal n-D array. More...
 
typedef Index_ Index
 Index type. More...
 
typedef LongIndex_ LongIndex
 Typically, strides in memory can be very large. More...
 
typedef Coord< kRank > TensorCoord
 Coordinate in logical tensor space. More...
 
typedef Coord< kStorageRank > StorageCoord
 Coordinate in storage n-D array. More...
 
typedef Coord< kStorageRank - 1 > StrideVector
 
typedef TensorRef< typename platform::remove_const< Storage >::type const, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ > ConstTensorRef
 Tensor reference to of constant value. More...
 
typedef TensorCoord Coord_t
 Coordinate in logical tensor space. More...
 
- - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + +

Public Member Functions

CUTLASS_HOST_DEVICE TensorView ()
 Default constructor. More...
 
CUTLASS_HOST_DEVICE TensorView (TensorRef_t const &_ref, Coord_t const &_size)
 Constructs a Tensor_view from a TensorRef and size. More...
 
CUTLASS_HOST_DEVICE bool good () const
 Returns true if the Tensor_view is bound to some memory. More...
 
CUTLASS_HOST_DEVICE T * data () const
 Returns a pointer to data. More...
 
CUTLASS_HOST_DEVICE void reset (TensorRef_t const &_ref=TensorRef_t(0), Coord_t const &_size=Coord_t())
 Updates the reference and size of a Tensor_view object. More...
 
CUTLASS_HOST_DEVICE TensorRef_t & ref ()
 Accesses the tensor reference pointing to data. More...
 
CUTLASS_HOST_DEVICE ConstTensorRef_t const_ref ()
 
CUTLASS_HOST_DEVICE TensorRef_t const & ref () const
 Accesses the tensor reference pointing to data. More...
 
CUTLASS_HOST_DEVICE Coord_t const & size () const
 Accesses the size. More...
 
CUTLASS_HOST_DEVICE int size (int dim) const
 Accesses the size. More...
 
CUTLASS_HOST_DEVICE Coord_t const & stride () const
 Accesses the stride. More...
 
CUTLASS_HOST_DEVICE int const & stride (int dim) const
 Accesses the stride. More...
 
CUTLASS_HOST_DEVICE TensorView & operator= (TensorView const &_tensor)
 Assigns the Tensor_view. More...
 
CUTLASS_HOST_DEVICE Offset_t offset (Coord_t const &coord) const
 Returns the index of an element. More...
 
CUTLASS_HOST_DEVICE bool contains (Coord_t const &coord) const
 Determines whether a location is within a tensor. More...
 
CUTLASS_HOST_DEVICE T & at (Coord_t const &coord) const
 Element-wise accessor. More...
 
T & operator[] (Coord< Rank > const &coord) const
 Element-wise accessor. More...
 
CUTLASS_HOST_DEVICE T & at (Offset_t idx) const
 Element-wise accessor. More...
 
CUTLASS_HOST_DEVICE TensorView< T > subview (Coord_t const &location, Coord_t size) const
 Returns a Tensor_view given location and size quantities. More...
 
- Public Member Functions inherited from cutlass::TensorRef< T, 4 >
CUTLASS_HOST_DEVICE TensorRef ()
 Default ctor. More...
 
CUTLASS_HOST_DEVICE TensorRef (Storage *ptr, Coord< Rank > stride)
 Constructs from a pointer, size, and stride. More...
 
CUTLASS_HOST_DEVICE void reset (Storage *ptr=nullptr, Coord< Rank > stride=Coord< Rank >(0))
 Updates the pointer, stride, and location within a TensorRef. More...
 
TensorRef< T, Rank > convert ()
 Conversion function. More...
 
CUTLASS_HOST_DEVICE bool good () const
 Returns true if the TensorRef may be safely accessed. More...
 
CUTLASS_HOST_DEVICE Storage * data () const
 Returns the pointer to referenced data. More...
 
CUTLASS_HOST_DEVICE Coord< Rank > const & stride () const
 Returns the stride of the tensor. More...
 
CUTLASS_HOST_DEVICE int const & stride (int dim) const
 Returns the stride of the tensor in the given dimension. More...
 
CUTLASS_HOST_DEVICE int leading_dim () const
 Returns the maximum stride element as the 'leading dimension'. More...
 
CUTLASS_HOST_DEVICE long long offset (Coord< Rank > const &coord) const
 Computes the offset of an index from the origin of the tensor. More...
 
CUTLASS_HOST_DEVICE Storage & at (Coord< Rank > const &coord) const
 Returns a reference to the element at a given Coord. More...
 
CUTLASS_HOST_DEVICE Storage & at (int idx) const
 Returns a reference to the element at a given Coord. More...
 
Storage & operator[] (Coord< Rank > const &coord) const
 Element-wise accessor. More...
 
Storage & operator[] (int idx) const
 Element-wise accessor. More...
 
CUTLASS_HOST_DEVICE TensorRef & advance (Coord< Rank > const &b)
 Adds an offset to the pointer. More...
 
CUTLASS_HOST_DEVICE TensorRef operator+ (Coord< Rank > const &b) const
 Returns a TensorRef offset by a given amount. More...
 
CUTLASS_HOST_DEVICE TensorRef operator- (Coord< Rank > const &b) const
 Returns a TensorRef offset by a given amount. More...
 
CUTLASS_HOST_DEVICE TensorView ()
 Default constructor. More...
 
CUTLASS_HOST_DEVICE TensorView (Base const &_ref, TensorCoord const &_size)
 Constructs a TensorView from a TensorRef and size. More...
 
CUTLASS_HOST_DEVICE TensorView (Storage *ptr, StrideVector const &stride, TensorCoord const &size)
 Constructs a TensorView from a pointer, a stride vector, and size. More...
 
CUTLASS_HOST_DEVICE TensorView (Storage *ptr, StorageCoord const &stride, TensorCoord const &size)
 Constructs a TensorView from a pointer, a stride vector, and size. More...
 
CUTLASS_HOST_DEVICE void reset (Base const &_ref=Base(), TensorCoord const &_size=TensorCoord())
 Updates the reference and size of a Tensor_view object. More...
 
CUTLASS_HOST_DEVICE TensorCoord const & size () const
 Accesses the size. More...
 
CUTLASS_HOST_DEVICE Index size (int dim) const
 Accesses the size. More...
 
CUTLASS_HOST_DEVICE TensorView & operator= (TensorView const &_tensor)
 Assigns the Tensor_view. More...
 
CUTLASS_HOST_DEVICE bool contains (TensorCoord const &coord) const
 Determines whether a location is within a tensor. More...
 
CUTLASS_HOST_DEVICE TensorRef ref () const
 Returns a TensorRef pointing to the first element of the tensor. More...
 
CUTLASS_HOST_DEVICE ConstTensorRef const_ref () const
 Returns a TensorRef pointing to the first element of the tensor. More...
 
CUTLASS_HOST_DEVICE TensorView subview (TensorCoord const &location, TensorCoord size) const
 Returns a Tensor_view given location and size quantities. More...
 
CUTLASS_HOST_DEVICE size_t capacity () const
 Returns the number of scalar elements needed to store tensor. More...
 
CUTLASS_HOST_DEVICE TensorView operator+ (TensorCoord const &b) const
 Returns a TensorView offset by a given amount. More...
 
CUTLASS_HOST_DEVICE TensorView & operator+= (TensorCoord const &b)
 Returns a TensorRef offset by a given amount. More...
 
CUTLASS_HOST_DEVICE TensorView operator- (TensorCoord const &b) const
 Returns a TensorRef offset by a given amount. More...
 
CUTLASS_HOST_DEVICE TensorView & operator-= (TensorCoord const &b)
 Returns a TensorRef offset by a given amount. More...
 
- Public Member Functions inherited from cutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >
CUTLASS_HOST_DEVICE TensorRef (Storage *ptr=nullptr)
 Helper for 1-D memory. All higher ranks are projected onto the fastest changing rank. More...
 
CUTLASS_HOST_DEVICE TensorRef (Storage *ptr, Index ldm)
 Helper to construct from a pointer and single stride element for 2-D pitch linear memory. More...
 
CUTLASS_HOST_DEVICE TensorRef (Storage *ptr, StrideVector const &stride)
 Constructs from a single pointer and stride vector. More...
 
CUTLASS_HOST_DEVICE TensorRef (Storage *ptr, StorageCoord const &stride)
 
CUTLASS_HOST_DEVICE TensorRef (TensorRef< typename platform::remove_const< Storage >::type, kRank, MapFunc, kStorageRank, Index, LongIndex > const &ref)
 Enables conversion from TensorRef of non-const type. More...
 
CUTLASS_HOST_DEVICE ConstTensorRef const_ref () const
 Returns a reference to constant-valued tensor. More...
 
CUTLASS_HOST_DEVICE void reset (Storage *ptr=nullptr)
 Updates only the pointer. More...
 
CUTLASS_HOST_DEVICE void reset (Storage *ptr, StorageCoord const &stride)
 Updates the pointer, stride, and location within a TensorRef. More...
 
CUTLASS_HOST_DEVICE bool good () const
 Returns true if the TensorRef may be safely accessed. More...
 
CUTLASS_HOST_DEVICE Storage * data () const
 Returns the pointer to referenced data. More...
 
CUTLASS_HOST_DEVICE StorageCoord stride () const
 Returns the stride of the tensor. More...
 
CUTLASS_HOST_DEVICE Index stride (int dim) const
 Returns the stride of the tensor in the given dimension. More...
 
CUTLASS_HOST_DEVICE Index leading_dim (int idx=0) const
 Returns the maximum stride element as the 'leading dimension'. More...
 
CUTLASS_HOST_DEVICE StorageCoord map (TensorCoord const &coord) const
 Maps a logical coordinate to an n-D array in memory. More...
 
CUTLASS_HOST_DEVICE LongIndex offset (TensorCoord const &coord) const
 Computes the offset of an index from the origin of the tensor. More...
 
CUTLASS_HOST_DEVICE Storage & at (TensorCoord const &coord) const
 Returns a reference to the element at a given Coord. More...
 
CUTLASS_HOST_DEVICE Storage & at (LongIndex idx) const
 Returns a reference to the element at a given linear index. More...
 
CUTLASS_HOST_DEVICE Storage & operator[] (TensorCoord const &coord) const
 Returns a reference to the element at a given Coord. More...
 
CUTLASS_HOST_DEVICE Storage & operator[] (LongIndex idx) const
 Returns a reference to the element at a given linear index. More...
 
CUTLASS_HOST_DEVICE TensorRef & add_pointer_offset (LongIndex delta)
 Adds an offset to each pointer. More...
 
CUTLASS_HOST_DEVICE TensorRef operator+ (TensorCoord const &b) const
 Returns a TensorRef offset by a given amount. More...
 
CUTLASS_HOST_DEVICE TensorRef & operator+= (TensorCoord const &b)
 Returns a TensorRef offset by a given amount. More...
 
CUTLASS_HOST_DEVICE TensorRef operator- (TensorCoord const &b) const
 Returns a TensorRef offset by a given amount. More...
 
CUTLASS_HOST_DEVICE TensorRef & operator-= (TensorCoord const &b)
 Returns a TensorRef offset by a given amount. More...
 
- - - - - - - + + + + + + + + + + + + +

Static Public Attributes

static int const Rank = TensorRef_t::Rank
 Rank of tensor. More...
 
- Static Public Attributes inherited from cutlass::TensorRef< T, 4 >
static int const Rank
 Rank of tensor. More...
 
static int const Rank = Base::kRank
 Logical rank of tensor index space. More...
 
- Static Public Attributes inherited from cutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >
static int const kRank = Rank_
 Logical rank of tensor index space. More...
 
static int const kStorageRank = StorageRank_
 Rank of internal storage. More...
 
static int const Rank = kRank
 Logical rank of tensor index space. More...
 

Member Typedef Documentation

- -

◆ Base

+ +

◆ Base

-template<typename T>
+template<typename Storage_ , int Rank_ = 4, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>
- +
typedef TensorRef<T, 4> cutlass::TensorView< T >::Basetypedef TensorRef<Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_> cutlass::TensorView< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::Base
- -

◆ ConstTensorRef_t

+ +

◆ ConstTensorRef

-template<typename T>
+template<typename Storage_ , int Rank_ = 4, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>
- +
typedef TensorRef<T const, 4> cutlass::TensorView< T >::ConstTensorRef_ttypedef TensorRef< typename platform::remove_const<Storage_>::type const, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_> cutlass::TensorView< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::ConstTensorRef
- -

◆ Coord_t

+ +

◆ ConstTensorRef_t

-template<typename T>
+template<typename Storage_ , int Rank_ = 4, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>
- +
typedef Coord<Rank> cutlass::TensorView< T >::Coord_ttypedef TensorRef::ConstTensorRef cutlass::TensorView< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::ConstTensorRef_t
- -

◆ Offset_t

+ +

◆ ConstTensorView

-template<typename T>
+template<typename Storage_ , int Rank_ = 4, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>
- +
typedef int cutlass::TensorView< T >::Offset_ttypedef TensorView< typename platform::remove_const<Storage>::type const, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_> cutlass::TensorView< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::ConstTensorView
- -

◆ TensorRef_t

+ +

◆ Coord_t

-template<typename T>
+template<typename Storage_ , int Rank_ = 4, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>
- + + +
typedef Base cutlass::TensorView< T >::TensorRef_ttypedef TensorCoord cutlass::TensorView< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::Coord_t
+
+ +
+ + +

◆ Index

+ +
+
+
+template<typename Storage_ , int Rank_ = 4, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>
+ + + + +
typedef Base::Index cutlass::TensorView< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::Index
+
+ +
+
+ +

◆ Offset_t

+ +
+
+
+template<typename Storage_ , int Rank_ = 4, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>
+ + + + +
typedef Base::LongIndex cutlass::TensorView< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::Offset_t
+
+ +
+
+ +

◆ Storage

+ +
+
+
+template<typename Storage_ , int Rank_ = 4, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>
+ + + + +
typedef Base::Storage cutlass::TensorView< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::Storage
+
+ +
+
+ +

◆ StorageCoord

+ +
+
+
+template<typename Storage_ , int Rank_ = 4, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>
+ + + + +
typedef TensorRef::StorageCoord cutlass::TensorView< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::StorageCoord
+
+ +
+
+ +

◆ StrideVector

+ +
+
+
+template<typename Storage_ , int Rank_ = 4, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>
+ + + + +
typedef TensorRef::StrideVector cutlass::TensorView< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::StrideVector
+
+

Stride vector in storage coordinate space Least significant stride is = 1 and not stored

+ +
+
+ +

◆ TensorCoord

+ +
+
+
+template<typename Storage_ , int Rank_ = 4, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>
+ + + + +
typedef TensorRef::TensorCoord cutlass::TensorView< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::TensorCoord
+
+ +
+
+ +

◆ TensorRef

+ +
+
+
+template<typename Storage_ , int Rank_ = 4, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>
+ + + + +
typedef Base cutlass::TensorView< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::TensorRef
+
+ +
+
+ +

◆ TensorRef_t

+ +
+
+
+template<typename Storage_ , int Rank_ = 4, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>
+ + +
typedef TensorRef cutlass::TensorView< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::TensorRef_t
@@ -321,19 +517,19 @@ template<typename T>

Constructor & Destructor Documentation

- -

◆ TensorView() [1/2]

+ +

◆ TensorView() [1/4]

-template<typename T>
+template<typename Storage_ , int Rank_ = 4, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>
+ + + + + + + + + + + + + + + @@ -109,6 +124,9 @@ Files + + + @@ -124,9 +142,6 @@ Files - - - @@ -145,12 +160,21 @@ Files + + + + + + + + + @@ -170,7 +194,7 @@ Files
-Generated on Thu May 3 2018 16:36:15 for Cutlass by   +Generated on Wed Sep 19 2018 13:58:53 for Cutlass by   doxygen 1.8.14
diff --git a/docs/dir_c5917a9a879e9a6c73eaf5237444ab84.html b/docs/dir_c5917a9a879e9a6c73eaf5237444ab84.html index a66eb22f..9011cf40 100644 --- a/docs/dir_c5917a9a879e9a6c73eaf5237444ab84.html +++ b/docs/dir_c5917a9a879e9a6c73eaf5237444ab84.html @@ -79,12 +79,16 @@ $(function() {
- + @@ -348,27 +544,27 @@ template<typename T> - -

◆ TensorView() [2/2]

+ +

◆ TensorView() [2/4]

-template<typename T>
+template<typename Storage_ , int Rank_ = 4, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>
CUTLASS_HOST_DEVICE cutlass::TensorView< T >::TensorView CUTLASS_HOST_DEVICE cutlass::TensorView< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::TensorView ( )
-Generated on Thu May 3 2018 16:36:12 for Cutlass by   +Generated on Wed Sep 19 2018 13:58:50 for Cutlass by   doxygen 1.8.14
diff --git a/docs/debug_8h_source.html b/docs/debug_8h_source.html index 881b4e3f..c404b411 100644 --- a/docs/debug_8h_source.html +++ b/docs/debug_8h_source.html @@ -81,7 +81,7 @@ $(function() {
-Generated on Thu May 3 2018 16:36:12 for Cutlass by   +Generated on Wed Sep 19 2018 13:58:50 for Cutlass by   doxygen 1.8.14
diff --git a/docs/dgemm__traits_8h.html b/docs/dgemm__traits_8h.html index eebc2f36..ac6d33b0 100644 --- a/docs/dgemm__traits_8h.html +++ b/docs/dgemm__traits_8h.html @@ -82,21 +82,21 @@ $(function() {

Defines structural traits of double-precision GEMM. More...

-
#include <cutlass/gemm/gemm.h>
-#include <cutlass/gemm/gemm_epilogue.h>
-#include <cutlass/gemm/gemm_epilogue_traits.h>
-#include <cutlass/gemm/gemm_global_tile.h>
-#include <cutlass/gemm/gemm_shared_tile.h>
-#include <cutlass/gemm/gemm_traits.h>
-#include <cutlass/gemm/thread_multiply_add.h>
+
#include "cutlass/gemm/gemm.h"
+#include "cutlass/gemm/gemm_epilogue.h"
+#include "cutlass/gemm/gemm_epilogue_traits.h"
+#include "cutlass/gemm/gemm_global_tile.h"
+#include "cutlass/gemm/gemm_shared_tile.h"
+#include "cutlass/gemm/gemm_traits.h"
+#include "cutlass/gemm/thread_multiply_add.h"

Go to the source code of this file.

- + - + - + @@ -386,332 +582,34 @@ template<typename T> -

Member Function Documentation

- -

◆ at() [1/2]

+ +

◆ TensorView() [3/4]

-template<typename T>
+template<typename Storage_ , int Rank_ = 4, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>
CUTLASS_HOST_DEVICE cutlass::TensorView< T >::TensorView CUTLASS_HOST_DEVICE cutlass::TensorView< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::TensorView (TensorRef_t const & Base const &  _ref,
Coord_t const & TensorCoord const &  _size 
- - -
- + - - - - -
CUTLASS_HOST_DEVICE T& cutlass::TensorView< T >::at CUTLASS_HOST_DEVICE cutlass::TensorView< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::TensorView (Coord_t const & coord) const
- 		-inline
-
- -
- - -

◆ at() [2/2]

- -
-
-
-template<typename T>
- - - - - -
- - - - - - - - -
CUTLASS_HOST_DEVICE T& cutlass::TensorView< T >::at (Offset_t idx) const
- 		-inline
-
- -
-
- -

◆ const_ref()

- -
-
-
-template<typename T>
- - - - - -
- - - - - - - -
CUTLASS_HOST_DEVICE ConstTensorRef_t cutlass::TensorView< T >::const_ref ()
- 		-inline
-
- -
-
- -

◆ contains()

- -
-
-
-template<typename T>
- - - - - -
- - - - - - - - -
CUTLASS_HOST_DEVICE bool cutlass::TensorView< T >::contains (Coord_t const & coord) const
- 		-inline
-
- -
-
- -

◆ data()

- -
-
-
-template<typename T>
- - - - - -
- - - - - - - -
CUTLASS_HOST_DEVICE T* cutlass::TensorView< T >::data () const
- 		-inline
-
- -
-
- -

◆ good()

- -
-
-
-template<typename T>
- - - - - -
- - - - - - - -
CUTLASS_HOST_DEVICE bool cutlass::TensorView< T >::good () const
- 		-inline
-
- -
-
- -

◆ offset()

- -
-
-
-template<typename T>
- - - - - -
- - - - - - - - -
CUTLASS_HOST_DEVICE Offset_t cutlass::TensorView< T >::offset (Coord_t const & coord) const
- 		-inline
-
- -
-
- -

◆ operator=()

- -
-
-
-template<typename T>
- - - - - -
- - - - - - - - -
CUTLASS_HOST_DEVICE TensorView& cutlass::TensorView< T >::operator= (TensorView< T > const & _tensor)
- 		-inline
-
- -
-
- -

◆ operator[]()

- -
-
-
-template<typename T>
- - - - - -
- - - - - - - - -
T& cutlass::TensorView< T >::operator[] (Coord< Rank > const & coord) const
- 		-inline
-
- -
-
- -

◆ ref() [1/2]

- -
-
-
-template<typename T>
- - - - - -
- - - - - - - -
CUTLASS_HOST_DEVICE TensorRef_t& cutlass::TensorView< T >::ref ()
- 		-inline
-
- -
-
- -

◆ ref() [2/2]

- -
-
-
-template<typename T>
- - - - - -
- - - - - - - -
CUTLASS_HOST_DEVICE TensorRef_t const& cutlass::TensorView< T >::ref () const
- 		-inline
-
- -
-
- -

◆ reset()

- -
-
-
-template<typename T>
- - -
- - - - - - + + - - + + + + + + + + @@ -728,19 +626,64 @@ template<typename T> - -

◆ size() [1/2]

+ +

◆ TensorView() [4/4]

-template<typename T>
+template<typename Storage_ , int Rank_ = 4, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>
CUTLASS_HOST_DEVICE void cutlass::TensorView< T >::reset (TensorRef_t const & _ref = TensorRef_t(0), Storage * ptr,
Coord_t const & _size = Coord_t() StrideVector const & stride,
TensorCoord const & size 
+ + +
- + + + + + + + + + + + + + + + + + + + + + + +
CUTLASS_HOST_DEVICE Coord_t const& cutlass::TensorView< T >::size CUTLASS_HOST_DEVICE cutlass::TensorView< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::TensorView (Storage * ptr,
StorageCoord const & stride,
TensorCoord const & size 
)
+ 		+inline
+
+ +
+ +

Member Function Documentation

+ +

◆ capacity()

+ +
+
+
+template<typename Storage_ , int Rank_ = 4, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>
+ + +
+ + + @@ -755,19 +698,306 @@ template<typename T> - -

◆ size() [2/2]

+ +

◆ const_ref()

-template<typename T>
+template<typename Storage_ , int Rank_ = 4, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>
CUTLASS_HOST_DEVICE size_t cutlass::TensorView< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::capacity ( ) const
+ + +
- + + + + + +
CUTLASS_HOST_DEVICE int cutlass::TensorView< T >::size CUTLASS_HOST_DEVICE ConstTensorRef cutlass::TensorView< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::const_ref () const
+ 		+inline
+
+ +
+ + +

◆ contains()

+ +
+
+
+template<typename Storage_ , int Rank_ = 4, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>
+ + + + + +
+ + + + + + + + +
CUTLASS_HOST_DEVICE bool cutlass::TensorView< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::contains (TensorCoord const & coord) const
+ 		+inline
+
+ +
+
+ +

◆ operator+()

+ +
+
+
+template<typename Storage_ , int Rank_ = 4, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>
+ + + + + +
+ + + + + + + + +
CUTLASS_HOST_DEVICE TensorView cutlass::TensorView< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::operator+ (TensorCoord const & b) const
+ 		+inline
+
+ +
+
+ +

◆ operator+=()

+ +
+
+
+template<typename Storage_ , int Rank_ = 4, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>
+ + + + + +
+ + + + + + + + +
CUTLASS_HOST_DEVICE TensorView& cutlass::TensorView< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::operator+= (TensorCoord const & b)
+ 		+inline
+
+ +
+
+ +

◆ operator-()

+ +
+
+
+template<typename Storage_ , int Rank_ = 4, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>
+ + + + + +
+ + + + + + + + +
CUTLASS_HOST_DEVICE TensorView cutlass::TensorView< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::operator- (TensorCoord const & b) const
+ 		+inline
+
+ +
+
+ +

◆ operator-=()

+ +
+
+
+template<typename Storage_ , int Rank_ = 4, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>
+ + + + + +
+ + + + + + + + +
CUTLASS_HOST_DEVICE TensorView& cutlass::TensorView< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::operator-= (TensorCoord const & b)
+ 		+inline
+
+ +
+
+ +

◆ operator=()

+ +
+
+
+template<typename Storage_ , int Rank_ = 4, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>
+ + + + + +
+ + + + + + + + +
CUTLASS_HOST_DEVICE TensorView& cutlass::TensorView< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::operator= (TensorView< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ > const & _tensor)
+ 		+inline
+
+ +
+
+ +

◆ ref()

+ +
+
+
+template<typename Storage_ , int Rank_ = 4, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>
+ + + + + +
+ + + + + + + +
CUTLASS_HOST_DEVICE TensorRef cutlass::TensorView< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::ref () const
+ 		+inline
+
+ +
+
+ +

◆ reset()

+ +
+
+
+template<typename Storage_ , int Rank_ = 4, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>
+ + + + + +
+ + + + + + + + + + + + + + + + + + +
CUTLASS_HOST_DEVICE void cutlass::TensorView< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::reset (Base const & _ref = Base(),
TensorCoord const & _size = TensorCoord() 
)
+ 		+inline
+
+ +
+
+ +

◆ size() [1/2]

+ +
+
+
+template<typename Storage_ , int Rank_ = 4, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>
+ + + + + +
+ + + + + + + +
CUTLASS_HOST_DEVICE TensorCoord const& cutlass::TensorView< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::size () const
+ 		+inline
+
+ +
+
+ +

◆ size() [2/2]

+ +
+
+
+template<typename Storage_ , int Rank_ = 4, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>
+ + +
+ + + @@ -783,82 +1013,27 @@ template<typename T> - -

◆ stride() [1/2]

+ +

◆ subview()

-template<typename T>
+template<typename Storage_ , int Rank_ = 4, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>
CUTLASS_HOST_DEVICE Index cutlass::TensorView< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::size ( int  dim)
- - -
- + - - - -
CUTLASS_HOST_DEVICE Coord_t const& cutlass::TensorView< T >::stride CUTLASS_HOST_DEVICE TensorView cutlass::TensorView< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::subview () const
- 		-inline
-
- -
- - -

◆ stride() [2/2]

- -
-
-
-template<typename T>
- - - - - -
- - - - - - - - -
CUTLASS_HOST_DEVICE int const& cutlass::TensorView< T >::stride (int dim) const
- 		-inline
-
- -
-
- -

◆ subview()

- -
-
-
-template<typename T>
- - -
- - - - - + - + @@ -877,19 +1052,19 @@ template<typename T>

Member Data Documentation

- -

◆ Rank

+ +

◆ Rank

-template<typename T>
+template<typename Storage_ , int Rank_ = 4, typename MapFunc_ = IdentityTensorMapFunc<Rank_>, int StorageRank_ = MapFunc_::kStorageRank, typename Index_ = int, typename LongIndex_ = long long>
CUTLASS_HOST_DEVICE TensorView<T> cutlass::TensorView< T >::subview (Coord_t const & TensorCoord const &  location,
Coord_t TensorCoord  size 
@@ -907,7 +1082,7 @@ template<typename T>
-Generated on Thu May 3 2018 16:36:13 for Cutlass by   +Generated on Wed Sep 19 2018 13:58:51 for Cutlass by   doxygen 1.8.14
diff --git a/docs/classcutlass_1_1TensorView.png b/docs/classcutlass_1_1TensorView.png index 40500e8a3a854639f0a6a25982cb66d99a343817..46861ac917aafe499ecdf03bd95fbfbe0af5725f 100644 GIT binary patch literal 1605 zcmZXUdpOg39LIkfvSm(}=_0wDHFlgFa%ZC=x1og`w_N5DvXsiC6B2u5WU5EBEWp^gQQzs`GuG@Av!uT)v;b-mfp&31=(6b>CJ10OYZD zRxSVlhDx*(L`rhfN_G!RuDwoe#vIpJAjU8ZIN^tQpgS?Pq0cyM$80w z3U^5!7n-hc;0d>3;d=FY7?gZZ0$|(E6+uRc0;0`DB8Anr;od`_e3fZ zgP6Y!!EQCF?U7pySwv-yzewO35{<6wexMAOo9@>t3`%ihnzSgE94}N9yk$8F`aHOo z-hbvtPy0X{4Pvr_Zp4yfJJeZiCO;^o68d^_^2I``LasCe8hu=V&PJxB2+gA~@RL z2esF~0cdO@_B&7%?|Ekg@K2L*5bC?Qc7*1reNAyyyFZ;vg1;ndq2k7IAFxS^WUW6%Qy z$dTZq2r7_ihXRmACS&aJv|}9`QW2@O3n}R0G;ufdmEhXM#@hN4;Y9wrFg51k zC;#?O89~$~%z2cy*COP}+PGl=fO+kW`DNYG6ZNMkd%A|BEK?o0IUpCCTZ@av zP@v)!+^RZ{g!k>gz&`8!MsGTq3im1^kk>BOJzEA3_@6KD?V*o7c!%Jrx2VcQ1X>95 z(FLr$xaKsc21W#h<-s#;>8(6Im;dyVC2Y3QKSy5fu)E0)ffCuMgX#W(&5FabQy?qg}2*e1I->6g(QfX{@h& kb_n-V=`AbG(io!+sToy*IYd+=FZuHUtToQ6$`YUOFIczx`2YX_ delta 645 zcmX@gvx!x)Gr-TCmrII^fq{Y7)59eQNG}5701jp#xwb6MdZMCLJ(HuSi(^Oy~60G%{^_8$>hx0QzUbI<8}*`%yia8DwU!glb&4V zO{(JcoTd5H-ZLue%d6#Q7M)3|YpoFWo@BLm&aCI9=Q3vI?N)rooBG#g|GZB!)o+`w zTkefkJgpct$MYV$^GwB)eZut{KWx4I#C`LutH!r1mP@r;91eRg*c;q_W1^Y!$r($` ztd3k~T0^0%nGUF0xDWj#~iq&-zia>omEbQ|hdvmdy3Oa9}>8}kb5 zkElO<*R!9&&`RpSdIy#dS^|uFxLg?4GYKA8jhp+=`^c1_ThS^Fb45EC;#id$!oPR; z|8Ck>Zzg@f>t?d-Yu_FDg6FR?idhLVPTTeMrNC#qO|O2hcyl)H_tn#L?&;otW#<3) z#@dJL1y3I-oBb!_TG>xSl?{spBKH)(7rXf8rp?za%1PoIB-6gW5^&$nw(-{Qts257 zws$-gxi7kF&Zpf+wu;8BT(@t1p=a2AdwJzsOWPdQ-fF(B8=VxrY2N4jO%Vse(hFl> zWIOv@eEt95C!JaE&d%7%&0Vs3vFCzX?ii~okNq3te}CURo$o*$l5by?D=eFG>Dn*m Y*ZfM>QKdJefGLZ?)78&qol`;+0NEli^Z)<= diff --git a/docs/classcutlass_1_1ZipTileIterator-members.html b/docs/classcutlass_1_1ZipTileIterator-members.html new file mode 100644 index 00000000..6de74a49 --- /dev/null +++ b/docs/classcutlass_1_1ZipTileIterator-members.html @@ -0,0 +1,125 @@ + + + + + + + +Cutlass: Member List + + + + + + + + + + +
+
+
- +
int const cutlass::TensorView< T >::Rank = TensorRef_t::Rankint const cutlass::TensorView< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::Rank = Base::kRank
+ + + + + +
+
Cutlass +
+
CUDA Templates for Linear Algebra Subroutines and Solvers
+
+ + + + + + + + + +
+
+ + +
+ +
+ + + +
+
+
cutlass::ZipTileIterator< First_, Second_ > Member List
+
+
+ +

This is the complete list of members for cutlass::ZipTileIterator< First_, Second_ >, including all inherited members.

+ + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + +
add_pointer_offset(Index offset)cutlass::ZipTileIterator< First_, Second_ >inline
decrement(int count=1)cutlass::ZipTileIterator< First_, Second_ >inline
First typedefcutlass::ZipTileIterator< First_, Second_ >
firstcutlass::ZipTileIterator< First_, Second_ >
Fragment typedefcutlass::ZipTileIterator< First_, Second_ >
increment(int count=1)cutlass::ZipTileIterator< First_, Second_ >inline
Index typedefcutlass::ZipTileIterator< First_, Second_ >
initialize_predicates(PredicateIterator predicate_it, Coord< 3 > const &bounds, Coord< 3 > const &block_offset=make_Coord(0, 0, 0))cutlass::ZipTileIterator< First_, Second_ >inline
initialize_predicates(PredicateIterator predicate_it, PredicateFunctor const &functor, Coord< 3 > const &block_offset)cutlass::ZipTileIterator< First_, Second_ >inline
load(Fragment &fragment) constcutlass::ZipTileIterator< First_, Second_ >inline
load(Fragment &fragment, Coord< 4 > const &offset) constcutlass::ZipTileIterator< First_, Second_ >inline
load(Fragment &fragment, PredicateIterator pred_it) constcutlass::ZipTileIterator< First_, Second_ >inline
load_post_increment(Fragment &fragment)cutlass::ZipTileIterator< First_, Second_ >inline
load_post_increment(Fragment &fragment, Coord< 4 > const &offset)cutlass::ZipTileIterator< First_, Second_ >inline
load_post_increment(Fragment &fragment, PredicateIterator pred_it)cutlass::ZipTileIterator< First_, Second_ >inline
operator++()cutlass::ZipTileIterator< First_, Second_ >inline
operator+=(int count)cutlass::ZipTileIterator< First_, Second_ >inline
operator+=(Coord< 3 > const &offset)cutlass::ZipTileIterator< First_, Second_ >inline
operator--()cutlass::ZipTileIterator< First_, Second_ >inline
operator-=(int count)cutlass::ZipTileIterator< First_, Second_ >inline
PredicateVector typedefcutlass::ZipTileIterator< First_, Second_ >
secondcutlass::ZipTileIterator< First_, Second_ >
Second typedefcutlass::ZipTileIterator< First_, Second_ >
store(Fragment const &fragment) constcutlass::ZipTileIterator< First_, Second_ >inline
store(Fragment const &fragment, Coord< 4 > const &offset) constcutlass::ZipTileIterator< First_, Second_ >inline
store(Fragment const &fragment, PredicateIterator pred_it) constcutlass::ZipTileIterator< First_, Second_ >inline
store_post_increment(Fragment const &fragment)cutlass::ZipTileIterator< First_, Second_ >inline
store_post_increment(Fragment const &fragment, Coord< 4 > const &offset)cutlass::ZipTileIterator< First_, Second_ >inline
store_post_increment(Fragment const &fragment, PredicateIterator pred_it)cutlass::ZipTileIterator< First_, Second_ >inline
TensorRef typedefcutlass::ZipTileIterator< First_, Second_ >
ZipTileIterator()cutlass::ZipTileIterator< First_, Second_ >inline
ZipTileIterator(Params const &_params, Coord< 3 > const &threadblock_offset=make_Coord(0, 0, 0))cutlass::ZipTileIterator< First_, Second_ >inline
ZipTileIterator(First const &_first, Second const &_second)cutlass::ZipTileIterator< First_, Second_ >inline
ZipTileIterator(TensorRef const &ref)cutlass::ZipTileIterator< First_, Second_ >inline
ZipTileIterator(Params const &_params, TensorRef const &ref)cutlass::ZipTileIterator< First_, Second_ >inline
+ +
+Generated on Wed Sep 19 2018 13:58:51 for Cutlass by   +doxygen + 1.8.14 +
+ + diff --git a/docs/classcutlass_1_1ZipTileIterator.html b/docs/classcutlass_1_1ZipTileIterator.html new file mode 100644 index 00000000..7cf7a392 --- /dev/null +++ b/docs/classcutlass_1_1ZipTileIterator.html @@ -0,0 +1,1290 @@ + + + + + + + +Cutlass: cutlass::ZipTileIterator< First_, Second_ > Class Template Reference + + + + + + + + + + +
+
+ + + + + + +
+
Cutlass +
+
CUDA Templates for Linear Algebra Subroutines and Solvers
+
+
+ + + + + + + + +
+
+ + +
+ +
+ + +
+
+
+Classes | +Public Types | +Public Member Functions | +Public Attributes | +List of all members
+
+
cutlass::ZipTileIterator< First_, Second_ > Class Template Reference
+
+
+ +

Constructs an iterator from a pair of iterators. +

+ +

#include <zip_tile_iterator.h>

+ + + + + +

+Classes

struct  Params
 Params object. More...
 
+ + + + + + + + + + + + + + + + + + + +

+Public Types

typedef First_ First
 First iterator type. More...
 
typedef Second_ Second
 Second iterator type. More...
 
typedef ZipFragment< typename First::Fragment, typename Second::Fragment > Fragment
 Fragment type. More...
 
typedef First::PredicateVector PredicateVector
 Predicate vector. More...
 
typedef First::Index Index
 Index type. More...
 
typedef ZipTensorRef< typename First::TensorRef, typename Second::TensorRef > TensorRef
 Tensor reference. More...
 
+ + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + +

+Public Member Functions

CUTLASS_DEVICE ZipTileIterator ()
 Default constructor. More...
 
CUTLASS_DEVICE ZipTileIterator (Params const &_params, Coord< 3 > const &threadblock_offset=make_Coord(0, 0, 0))
 Constructs a zip iterator from params. More...
 
CUTLASS_DEVICE ZipTileIterator (First const &_first, Second const &_second)
 Constructs a zip iterator from iterator instances. More...
 
CUTLASS_DEVICE ZipTileIterator (TensorRef const &ref)
 Constructs a zip iterator from iterator instances. More...
 
CUTLASS_DEVICE ZipTileIterator (Params const &_params, TensorRef const &ref)
 Constructs a zip iterator from iterator instances. More...
 
template<typename PredicateIterator >
CUTLASS_HOST_DEVICE void initialize_predicates (PredicateIterator predicate_it, Coord< 3 > const &bounds, Coord< 3 > const &block_offset=make_Coord(0, 0, 0))
 Initializes a predicate vector using a RegularTilePredicateFunctor. More...
 
template<typename PredicateIterator , typename PredicateFunctor >
CUTLASS_HOST_DEVICE void initialize_predicates (PredicateIterator predicate_it, PredicateFunctor const &functor, Coord< 3 > const &block_offset)
 Initializes a predicate vector using an arbitrary predicate functor. More...
 
template<typename Fragment >
CUTLASS_DEVICE void load_post_increment (Fragment &fragment)
 Loads a fragment and increments without predicates. More...
 
template<typename Fragment >
CUTLASS_DEVICE void load_post_increment (Fragment &fragment, Coord< 4 > const &offset)
 Loads a fragment and increments without predicates. More...
 
template<typename Fragment >
CUTLASS_DEVICE void load (Fragment &fragment) const
 Loads a fragment without predicates. More...
 
template<typename Fragment >
CUTLASS_DEVICE void load (Fragment &fragment, Coord< 4 > const &offset) const
 Loads a fragment without predicates. More...
 
template<typename Fragment >
CUTLASS_DEVICE void store_post_increment (Fragment const &fragment)
 Stores a fragment and increments without predicates. More...
 
template<typename Fragment >
CUTLASS_DEVICE void store_post_increment (Fragment const &fragment, Coord< 4 > const &offset)
 Stores a fragment and increments without predicates. More...
 
template<typename Fragment >
CUTLASS_DEVICE void store (Fragment const &fragment) const
 Stores a fragment without predicates. More...
 
template<typename Fragment >
CUTLASS_DEVICE void store (Fragment const &fragment, Coord< 4 > const &offset) const
 Stores a fragment without predicates. More...
 
template<typename Fragment , typename PredicateIterator >
CUTLASS_DEVICE void load_post_increment (Fragment &fragment, PredicateIterator pred_it)
 Loads a fragment and increments, using predicates. More...
 
template<typename Fragment , typename PredicateIterator >
CUTLASS_DEVICE void load (Fragment &fragment, PredicateIterator pred_it) const
 Loads a fragment with predicates. More...
 
template<typename Fragment , typename PredicateIterator >
CUTLASS_DEVICE void store_post_increment (Fragment const &fragment, PredicateIterator pred_it)
 Loads a fragment and increments, using predicates. More...
 
template<typename Fragment , typename PredicateIterator >
CUTLASS_DEVICE void store (Fragment const &fragment, PredicateIterator pred_it) const
 Loads a fragment with predicates. More...
 
CUTLASS_DEVICE ZipTileIterator & increment (int count=1)
 Increments store iterator to next tile. More...
 
CUTLASS_DEVICE ZipTileIterator & operator++ ()
 Increments to next tile. More...
 
CUTLASS_DEVICE ZipTileIterator & operator+= (int count)
 
CUTLASS_DEVICE ZipTileIterator & operator+= (Coord< 3 > const &offset)
 Adds a vector offset to the underlying iterators. More...
 
CUTLASS_DEVICE ZipTileIterator & decrement (int count=1)
 Increments store iterator to previous tile. More...
 
CUTLASS_DEVICE ZipTileIterator & operator-- ()
 Increments to subsequent tile. More...
 
CUTLASS_DEVICE ZipTileIterator & operator-= (int count)
 Decrements to previous tile. More...
 
CUTLASS_DEVICE void add_pointer_offset (Index offset)
 Adds an offset to both iterators. More...
 
+ + + + + + + +

+Public Attributes

First first
 First iterator. More...
 
Second second
 Second iterator. More...
 
+

Member Typedef Documentation

+ +

◆ First

+ +
+
+
+template<typename First_ , typename Second_ >
+ + + + +
typedef First_ cutlass::ZipTileIterator< First_, Second_ >::First
+
+ +
+
+ +

◆ Fragment

+ +
+
+
+template<typename First_ , typename Second_ >
+ + + + +
typedef ZipFragment<typename First::Fragment, typename Second::Fragment> cutlass::ZipTileIterator< First_, Second_ >::Fragment
+
+ +
+
+ +

◆ Index

+ +
+
+
+template<typename First_ , typename Second_ >
+ + + + +
typedef First::Index cutlass::ZipTileIterator< First_, Second_ >::Index
+
+ +
+
+ +

◆ PredicateVector

+ +
+
+
+template<typename First_ , typename Second_ >
+ + + + +
typedef First::PredicateVector cutlass::ZipTileIterator< First_, Second_ >::PredicateVector
+
+ +
+
+ +

◆ Second

+ +
+
+
+template<typename First_ , typename Second_ >
+ + + + +
typedef Second_ cutlass::ZipTileIterator< First_, Second_ >::Second
+
+ +
+
+ +

◆ TensorRef

+ +
+
+
+template<typename First_ , typename Second_ >
+ + + + +
typedef ZipTensorRef< typename First::TensorRef, typename Second::TensorRef> cutlass::ZipTileIterator< First_, Second_ >::TensorRef
+
+ +
+
+

Constructor & Destructor Documentation

+ +

◆ ZipTileIterator() [1/5]

+ +
+
+
+template<typename First_ , typename Second_ >
+ + + + + +
+ + + + + + + +
CUTLASS_DEVICE cutlass::ZipTileIterator< First_, Second_ >::ZipTileIterator ()
+ 		+inline
+
+ +
+
+ +

◆ ZipTileIterator() [2/5]

+ +
+
+
+template<typename First_ , typename Second_ >
+ + + + + +
+ + + + + + + + + + + + + + + + + + +
CUTLASS_DEVICE cutlass::ZipTileIterator< First_, Second_ >::ZipTileIterator (Params const & _params,
Coord< 3 > const & threadblock_offset = make_Coord(0, 0, 0) 
)
+ 		+inline
+
+ +
+
+ +

◆ ZipTileIterator() [3/5]

+ +
+
+
+template<typename First_ , typename Second_ >
+ + + + + +
+ + + + + + + + + + + + + + + + + + +
CUTLASS_DEVICE cutlass::ZipTileIterator< First_, Second_ >::ZipTileIterator (First const & _first,
Second const & _second 
)
+ 		+inline
+
+ +
+
+ +

◆ ZipTileIterator() [4/5]

+ +
+
+
+template<typename First_ , typename Second_ >
+ + + + + +
+ + + + + + + + +
CUTLASS_DEVICE cutlass::ZipTileIterator< First_, Second_ >::ZipTileIterator (TensorRef const & ref)
+ 		+inline
+
+ +
+
+ +

◆ ZipTileIterator() [5/5]

+ +
+
+
+template<typename First_ , typename Second_ >
+ + + + + +
+ + + + + + + + + + + + + + + + + + +
CUTLASS_DEVICE cutlass::ZipTileIterator< First_, Second_ >::ZipTileIterator (Params const & _params,
TensorRef const & ref 
)
+ 		+inline
+
+ +
+
+

Member Function Documentation

+ +

◆ add_pointer_offset()

+ +
+
+
+template<typename First_ , typename Second_ >
+ + + + + +
+ + + + + + + + +
CUTLASS_DEVICE void cutlass::ZipTileIterator< First_, Second_ >::add_pointer_offset (Index offset)
+ 		+inline
+
+ +
+
+ +

◆ decrement()

+ +
+
+
+template<typename First_ , typename Second_ >
+ + + + + +
+ + + + + + + + +
CUTLASS_DEVICE ZipTileIterator& cutlass::ZipTileIterator< First_, Second_ >::decrement (int count = 1)
+ 		+inline
+
+ +
+
+ +

◆ increment()

+ +
+
+
+template<typename First_ , typename Second_ >
+ + + + + +
+ + + + + + + + +
CUTLASS_DEVICE ZipTileIterator& cutlass::ZipTileIterator< First_, Second_ >::increment (int count = 1)
+ 		+inline
+
+ +
+
+ +

◆ initialize_predicates() [1/2]

+ +
+
+
+template<typename First_ , typename Second_ >
+
+template<typename PredicateIterator >
+ + + + + +
+ + + + + + + + + + + + + + + + + + + + + + + + +
CUTLASS_HOST_DEVICE void cutlass::ZipTileIterator< First_, Second_ >::initialize_predicates (PredicateIterator predicate_it,
Coord< 3 > const & bounds,
Coord< 3 > const & block_offset = make_Coord(0,                                                                                           0,                                                                                           0) 
)
+ 		+inline
+
+ +
+
+ +

◆ initialize_predicates() [2/2]

+ +
+
+
+template<typename First_ , typename Second_ >
+
+template<typename PredicateIterator , typename PredicateFunctor >
+ + + + + +
+ + + + + + + + + + + + + + + + + + + + + + + + +
CUTLASS_HOST_DEVICE void cutlass::ZipTileIterator< First_, Second_ >::initialize_predicates (PredicateIterator predicate_it,
PredicateFunctor const & functor,
Coord< 3 > const & block_offset 
)
+ 		+inline
+
+ +
+
+ +

◆ load() [1/3]

+ +
+
+
+template<typename First_ , typename Second_ >
+
+template<typename Fragment >
+ + + + + +
+ + + + + + + + +
CUTLASS_DEVICE void cutlass::ZipTileIterator< First_, Second_ >::load (Fragment & fragment) const
+ 		+inline
+
+ +
+
+ +

◆ load() [2/3]

+ +
+
+
+template<typename First_ , typename Second_ >
+
+template<typename Fragment >
+ + + + + +
+ + + + + + + + + + + + + + + + + + +
CUTLASS_DEVICE void cutlass::ZipTileIterator< First_, Second_ >::load (Fragment & fragment,
Coord< 4 > const & offset 
) const
+ 		+inline
+
+ +
+
+ +

◆ load() [3/3]

+ +
+
+
+template<typename First_ , typename Second_ >
+
+template<typename Fragment , typename PredicateIterator >
+ + + + + +
+ + + + + + + + + + + + + + + + + + +
CUTLASS_DEVICE void cutlass::ZipTileIterator< First_, Second_ >::load (Fragment & fragment,
PredicateIterator pred_it 
) const
+ 		+inline
+
+ +
+
+ +

◆ load_post_increment() [1/3]

+ +
+
+
+template<typename First_ , typename Second_ >
+
+template<typename Fragment >
+ + + + + +
+ + + + + + + + +
CUTLASS_DEVICE void cutlass::ZipTileIterator< First_, Second_ >::load_post_increment (Fragment & fragment)
+ 		+inline
+
+ +
+
+ +

◆ load_post_increment() [2/3]

+ +
+
+
+template<typename First_ , typename Second_ >
+
+template<typename Fragment >
+ + + + + +
+ + + + + + + + + + + + + + + + + + +
CUTLASS_DEVICE void cutlass::ZipTileIterator< First_, Second_ >::load_post_increment (Fragment & fragment,
Coord< 4 > const & offset 
)
+ 		+inline
+
+ +
+
+ +

◆ load_post_increment() [3/3]

+ +
+
+
+template<typename First_ , typename Second_ >
+
+template<typename Fragment , typename PredicateIterator >
+ + + + + +
+ + + + + + + + + + + + + + + + + + +
CUTLASS_DEVICE void cutlass::ZipTileIterator< First_, Second_ >::load_post_increment (Fragment & fragment,
PredicateIterator pred_it 
)
+ 		+inline
+
+ +
+
+ +

◆ operator++()

+ +
+
+
+template<typename First_ , typename Second_ >
+ + + + + +
+ + + + + + + +
CUTLASS_DEVICE ZipTileIterator& cutlass::ZipTileIterator< First_, Second_ >::operator++ ()
+ 		+inline
+
+ +
+
+ +

◆ operator+=() [1/2]

+ +
+
+
+template<typename First_ , typename Second_ >
+ + + + + +
+ + + + + + + + +
CUTLASS_DEVICE ZipTileIterator& cutlass::ZipTileIterator< First_, Second_ >::operator+= (int count)
+ 		+inline
+
+ +
+
+ +

◆ operator+=() [2/2]

+ +
+
+
+template<typename First_ , typename Second_ >
+ + + + + +
+ + + + + + + + +
CUTLASS_DEVICE ZipTileIterator& cutlass::ZipTileIterator< First_, Second_ >::operator+= (Coord< 3 > const & offset)
+ 		+inline
+
+ +
+
+ +

◆ operator--()

+ +
+
+
+template<typename First_ , typename Second_ >
+ + + + + +
+ + + + + + + +
CUTLASS_DEVICE ZipTileIterator& cutlass::ZipTileIterator< First_, Second_ >::operator-- ()
+ 		+inline
+
+ +
+
+ +

◆ operator-=()

+ +
+
+
+template<typename First_ , typename Second_ >
+ + + + + +
+ + + + + + + + +
CUTLASS_DEVICE ZipTileIterator& cutlass::ZipTileIterator< First_, Second_ >::operator-= (int count)
+ 		+inline
+
+ +
+
+ +

◆ store() [1/3]

+ +
+
+
+template<typename First_ , typename Second_ >
+
+template<typename Fragment >
+ + + + + +
+ + + + + + + + +
CUTLASS_DEVICE void cutlass::ZipTileIterator< First_, Second_ >::store (Fragment const & fragment) const
+ 		+inline
+
+ +
+
+ +

◆ store() [2/3]

+ +
+
+
+template<typename First_ , typename Second_ >
+
+template<typename Fragment >
+ + + + + +
+ + + + + + + + + + + + + + + + + + +
CUTLASS_DEVICE void cutlass::ZipTileIterator< First_, Second_ >::store (Fragment const & fragment,
Coord< 4 > const & offset 
) const
+ 		+inline
+
+ +
+
+ +

◆ store() [3/3]

+ +
+
+
+template<typename First_ , typename Second_ >
+
+template<typename Fragment , typename PredicateIterator >
+ + + + + +
+ + + + + + + + + + + + + + + + + + +
CUTLASS_DEVICE void cutlass::ZipTileIterator< First_, Second_ >::store (Fragment const & fragment,
PredicateIterator pred_it 
) const
+ 		+inline
+
+ +
+
+ +

◆ store_post_increment() [1/3]

+ +
+
+
+template<typename First_ , typename Second_ >
+
+template<typename Fragment >
+ + + + + +
+ + + + + + + + +
CUTLASS_DEVICE void cutlass::ZipTileIterator< First_, Second_ >::store_post_increment (Fragment const & fragment)
+ 		+inline
+
+ +
+
+ +

◆ store_post_increment() [2/3]

+ +
+
+
+template<typename First_ , typename Second_ >
+
+template<typename Fragment >
+ + + + + +
+ + + + + + + + + + + + + + + + + + +
CUTLASS_DEVICE void cutlass::ZipTileIterator< First_, Second_ >::store_post_increment (Fragment const & fragment,
Coord< 4 > const & offset 
)
+ 		+inline
+
+ +
+
+ +

◆ store_post_increment() [3/3]

+ +
+
+
+template<typename First_ , typename Second_ >
+
+template<typename Fragment , typename PredicateIterator >
+ + + + + +
+ + + + + + + + + + + + + + + + + + +
CUTLASS_DEVICE void cutlass::ZipTileIterator< First_, Second_ >::store_post_increment (Fragment const & fragment,
PredicateIterator pred_it 
)
+ 		+inline
+
+ +
+
+

Member Data Documentation

+ +

◆ first

+ +
+
+
+template<typename First_ , typename Second_ >
+ + + + +
First cutlass::ZipTileIterator< First_, Second_ >::first
+
+ +
+
+ +

◆ second

+ +
+
+
+template<typename First_ , typename Second_ >
+ + + + +
Second cutlass::ZipTileIterator< First_, Second_ >::second
+
+ +
+
+
The documentation for this class was generated from the following file: +
+ +
+Generated on Wed Sep 19 2018 13:58:51 for Cutlass by   +doxygen + 1.8.14 +
+ + diff --git a/docs/classcutlass_1_1detail_1_1ScalarOrPointer-members.html b/docs/classcutlass_1_1detail_1_1ScalarOrPointer-members.html new file mode 100644 index 00000000..8da71466 --- /dev/null +++ b/docs/classcutlass_1_1detail_1_1ScalarOrPointer-members.html @@ -0,0 +1,101 @@ + + + + + + + +Cutlass: Member List + + + + + + + + + + +
+
+ + + + + + +
+
Cutlass +
+
CUDA Templates for Linear Algebra Subroutines and Solvers
+
+
+ + + + + + + + +
+
+ + +
+ +
+ + +
+
+
+
cutlass::detail::ScalarOrPointer< Scalar_ > Member List
+
+
+ +

This is the complete list of members for cutlass::detail::ScalarOrPointer< Scalar_ >, including all inherited members.

+ + + + + + + + + + + + +
get() constcutlass::detail::ScalarOrPointer< Scalar_ >inline
get_ptr() constcutlass::detail::ScalarOrPointer< Scalar_ >inline
get_scalar() constcutlass::detail::ScalarOrPointer< Scalar_ >inline
is_pointer() constcutlass::detail::ScalarOrPointer< Scalar_ >inline
operator Scalar() constcutlass::detail::ScalarOrPointer< Scalar_ >inline
operator=(Scalar const &scalar_)cutlass::detail::ScalarOrPointer< Scalar_ >inline
operator=(Scalar const *ptr_)cutlass::detail::ScalarOrPointer< Scalar_ >inline
Scalar typedefcutlass::detail::ScalarOrPointer< Scalar_ >
ScalarOrPointer()cutlass::detail::ScalarOrPointer< Scalar_ >inline
ScalarOrPointer(Scalar const &val)cutlass::detail::ScalarOrPointer< Scalar_ >inline
ScalarOrPointer(Scalar const *ptr_)cutlass::detail::ScalarOrPointer< Scalar_ >inline
+ +
+Generated on Wed Sep 19 2018 13:58:51 for Cutlass by   +doxygen + 1.8.14 +
+ + diff --git a/docs/classcutlass_1_1detail_1_1ScalarOrPointer.html b/docs/classcutlass_1_1detail_1_1ScalarOrPointer.html new file mode 100644 index 00000000..6a28c38f --- /dev/null +++ b/docs/classcutlass_1_1detail_1_1ScalarOrPointer.html @@ -0,0 +1,434 @@ + + + + + + + +Cutlass: cutlass::detail::ScalarOrPointer< Scalar_ > Class Template Reference + + + + + + + + + + +
+
+ + + + + + +
+
Cutlass +
+
CUDA Templates for Linear Algebra Subroutines and Solvers
+
+
+ + + + + + + + +
+
+ + +
+ +
+ + +
+
+
+Public Types | +Public Member Functions | +List of all members
+
+
cutlass::detail::ScalarOrPointer< Scalar_ > Class Template Reference
+
+
+ +

#include <scalar_or_pointer.h>

+ + + + + +

+Public Types

typedef Scalar_ Scalar
 Underlying scalar type. More...
 
+ + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + +

+Public Member Functions

CUTLASS_HOST_DEVICE ScalarOrPointer ()
 Default ctor. More...
 
CUTLASS_HOST_DEVICE ScalarOrPointer (Scalar const &val)
 Object behaves as a scalar. More...
 
CUTLASS_HOST_DEVICE ScalarOrPointer (Scalar const *ptr_)
 Object behaves as a scalar. More...
 
CUTLASS_HOST_DEVICE bool is_pointer () const
 Returns true if is pointer. More...
 
CUTLASS_HOST_DEVICE Scalar const * get_ptr () const
 Gets the pointer value. More...
 
CUTLASS_HOST_DEVICE Scalar get_scalar () const
 Gets the pointer value. More...
 
CUTLASS_HOST_DEVICE ScalarOrPointer & operator= (Scalar const &scalar_)
 Assigns to a scalar and sets pointer to nullptr. More...
 
CUTLASS_HOST_DEVICE ScalarOrPointer & operator= (Scalar const *ptr_)
 Assigns to a pointer value. More...
 
CUTLASS_HOST_DEVICE Scalar get () const
 Access the element. More...
 
CUTLASS_HOST_DEVICE operator Scalar () const
 Accesses the element. More...
 
+

Detailed Description

+

template<typename Scalar_>
+class cutlass::detail::ScalarOrPointer< Scalar_ >

+ +

Helper class defines an object which operates as either a scalar or a pointer. If the pointer is non-null, it is dereferenced when the object is accessed.

+

Member Typedef Documentation

+ +

◆ Scalar

+ +
+
+
+template<typename Scalar_>
+ + + + +
typedef Scalar_ cutlass::detail::ScalarOrPointer< Scalar_ >::Scalar
+
+ +
+
+

Constructor & Destructor Documentation

+ +

◆ ScalarOrPointer() [1/3]

+ +
+
+
+template<typename Scalar_>
+ + + + + +
+ + + + + + + +
CUTLASS_HOST_DEVICE cutlass::detail::ScalarOrPointer< Scalar_ >::ScalarOrPointer ()
+ 		+inline
+
+ +
+
+ +

◆ ScalarOrPointer() [2/3]

+ +
+
+
+template<typename Scalar_>
+ + + + + +
+ + + + + + + + +
CUTLASS_HOST_DEVICE cutlass::detail::ScalarOrPointer< Scalar_ >::ScalarOrPointer (Scalar const & val)
+ 		+inline
+
+ +
+
+ +

◆ ScalarOrPointer() [3/3]

+ +
+
+
+template<typename Scalar_>
+ + + + + +
+ + + + + + + + +
CUTLASS_HOST_DEVICE cutlass::detail::ScalarOrPointer< Scalar_ >::ScalarOrPointer (Scalar const * ptr_)
+ 		+inline
+
+ +
+
+

Member Function Documentation

+ +

◆ get()

+ +
+
+
+template<typename Scalar_>
+ + + + + +
+ + + + + + + +
CUTLASS_HOST_DEVICE Scalar cutlass::detail::ScalarOrPointer< Scalar_ >::get () const
+ 		+inline
+
+ +
+
+ +

◆ get_ptr()

+ +
+
+
+template<typename Scalar_>
+ + + + + +
+ + + + + + + +
CUTLASS_HOST_DEVICE Scalar const* cutlass::detail::ScalarOrPointer< Scalar_ >::get_ptr () const
+ 		+inline
+
+ +
+
+ +

◆ get_scalar()

+ +
+
+
+template<typename Scalar_>
+ + + + + +
+ + + + + + + +
CUTLASS_HOST_DEVICE Scalar cutlass::detail::ScalarOrPointer< Scalar_ >::get_scalar () const
+ 		+inline
+
+ +
+
+ +

◆ is_pointer()

+ +
+
+
+template<typename Scalar_>
+ + + + + +
+ + + + + + + +
CUTLASS_HOST_DEVICE bool cutlass::detail::ScalarOrPointer< Scalar_ >::is_pointer () const
+ 		+inline
+
+ +
+
+ +

◆ operator Scalar()

+ +
+
+
+template<typename Scalar_>
+ + + + + +
+ + + + + + + +
CUTLASS_HOST_DEVICE cutlass::detail::ScalarOrPointer< Scalar_ >::operator Scalar () const
+ 		+inline
+
+ +
+
+ +

◆ operator=() [1/2]

+ +
+
+
+template<typename Scalar_>
+ + + + + +
+ + + + + + + + +
CUTLASS_HOST_DEVICE ScalarOrPointer& cutlass::detail::ScalarOrPointer< Scalar_ >::operator= (Scalar const & scalar_)
+ 		+inline
+
+ +
+
+ +

◆ operator=() [2/2]

+ +
+
+
+template<typename Scalar_>
+ + + + + +
+ + + + + + + + +
CUTLASS_HOST_DEVICE ScalarOrPointer& cutlass::detail::ScalarOrPointer< Scalar_ >::operator= (Scalar const * ptr_)
+ 		+inline
+
+ +
+
+
The documentation for this class was generated from the following file: +
+ +
+Generated on Wed Sep 19 2018 13:58:51 for Cutlass by   +doxygen + 1.8.14 +
+ + diff --git a/docs/classcutlass_1_1gemm_1_1LinearScalingDevicePtr_1_1Params-members.html b/docs/classcutlass_1_1gemm_1_1LinearScalingDevicePtr_1_1Params-members.html new file mode 100644 index 00000000..323b1406 --- /dev/null +++ b/docs/classcutlass_1_1gemm_1_1LinearScalingDevicePtr_1_1Params-members.html @@ -0,0 +1,98 @@ + + + + + + + +Cutlass: Member List + + + + + + + + + + +
+
+ + + + + + +
+
Cutlass +
+
CUDA Templates for Linear Algebra Subroutines and Solvers
+
+
+ + + + + + + + +
+
+ + +
+ +
+ + +
+
+
+
cutlass::gemm::LinearScalingDevicePtr< Scalar_, FragmentMultiplyAdd_ >::Params Member List
+
+
+ +

This is the complete list of members for cutlass::gemm::LinearScalingDevicePtr< Scalar_, FragmentMultiplyAdd_ >::Params, including all inherited members.

+ + + + + + + + + +
alpha() constcutlass::gemm::LinearScalingDevicePtr< Scalar_, FragmentMultiplyAdd_ >::Paramsinline
beta() constcutlass::gemm::LinearScalingDevicePtr< Scalar_, FragmentMultiplyAdd_ >::Paramsinline
initialize(Scalar alpha, Scalar beta)cutlass::gemm::LinearScalingDevicePtr< Scalar_, FragmentMultiplyAdd_ >::Paramsinline
initialize(Scalar const *alpha, Scalar const *beta)cutlass::gemm::LinearScalingDevicePtr< Scalar_, FragmentMultiplyAdd_ >::Paramsinline
initialize(GemmDesc_ const &desc)cutlass::gemm::LinearScalingDevicePtr< Scalar_, FragmentMultiplyAdd_ >::Paramsinline
Params()cutlass::gemm::LinearScalingDevicePtr< Scalar_, FragmentMultiplyAdd_ >::Paramsinline
Params(Scalar alpha, Scalar beta)cutlass::gemm::LinearScalingDevicePtr< Scalar_, FragmentMultiplyAdd_ >::Paramsinline
Params(Scalar const *alpha_ptr, Scalar const *beta_ptr)cutlass::gemm::LinearScalingDevicePtr< Scalar_, FragmentMultiplyAdd_ >::Paramsinline
+ +
+Generated on Wed Sep 19 2018 13:58:53 for Cutlass by   +doxygen + 1.8.14 +
+ + diff --git a/docs/classcutlass_1_1gemm_1_1LinearScalingDevicePtr_1_1Params.html b/docs/classcutlass_1_1gemm_1_1LinearScalingDevicePtr_1_1Params.html new file mode 100644 index 00000000..5fc5d05e --- /dev/null +++ b/docs/classcutlass_1_1gemm_1_1LinearScalingDevicePtr_1_1Params.html @@ -0,0 +1,389 @@ + + + + + + + +Cutlass: cutlass::gemm::LinearScalingDevicePtr< Scalar_, FragmentMultiplyAdd_ >::Params Class Reference + + + + + + + + + + +
+
+ + + + + + +
+
Cutlass +
+
CUDA Templates for Linear Algebra Subroutines and Solvers
+
+
+ + + + + + + + +
+
+ + +
+ +
+ + +
+
+
+Public Member Functions | +List of all members
+
+
cutlass::gemm::LinearScalingDevicePtr< Scalar_, FragmentMultiplyAdd_ >::Params Class Reference
+
+
+ +

The parameters. +

+ +

#include <linear_scaling_device_ptr.h>

+ + + + + + + + + + + + + + + + + + + + + + + + +

+Public Member Functions

CUTLASS_HOST_DEVICE Params ()
 
CUTLASS_HOST_DEVICE Params (Scalar alpha, Scalar beta)
 
CUTLASS_HOST_DEVICE Params (Scalar const *alpha_ptr, Scalar const *beta_ptr)
 
CUTLASS_HOST_DEVICE int initialize (Scalar alpha, Scalar beta)
 Initialize the parameters. More...
 
CUTLASS_HOST_DEVICE int initialize (Scalar const *alpha, Scalar const *beta)
 Initialize the parameters. More...
 
template<typename GemmDesc_ >
CUTLASS_HOST_DEVICE int initialize (GemmDesc_ const &desc)
 Initialize the parameters. More...
 
CUTLASS_HOST_DEVICE Scalar alpha () const
 Gets the alpha scalar. More...
 
CUTLASS_HOST_DEVICE Scalar beta () const
 Gets the beta scalar. More...
 
+

Constructor & Destructor Documentation

+ +

◆ Params() [1/3]

+ +
+
+
+template<typename Scalar_ , typename FragmentMultiplyAdd_ = FragmentMultiplyAdd<Scalar_, Scalar_>>
+ + + + + +
+ + + + + + + +
CUTLASS_HOST_DEVICE cutlass::gemm::LinearScalingDevicePtr< Scalar_, FragmentMultiplyAdd_ >::Params::Params ()
+ 		+inline
+
+ +
+
+ +

◆ Params() [2/3]

+ +
+
+
+template<typename Scalar_ , typename FragmentMultiplyAdd_ = FragmentMultiplyAdd<Scalar_, Scalar_>>
+ + + + + +
+ + + + + + + + + + + + + + + + + + +
CUTLASS_HOST_DEVICE cutlass::gemm::LinearScalingDevicePtr< Scalar_, FragmentMultiplyAdd_ >::Params::Params (Scalar alpha,
Scalar beta 
)
+ 		+inline
+
+ +
+
+ +

◆ Params() [3/3]

+ +
+
+
+template<typename Scalar_ , typename FragmentMultiplyAdd_ = FragmentMultiplyAdd<Scalar_, Scalar_>>
+ + + + + +
+ + + + + + + + + + + + + + + + + + +
CUTLASS_HOST_DEVICE cutlass::gemm::LinearScalingDevicePtr< Scalar_, FragmentMultiplyAdd_ >::Params::Params (Scalar const * alpha_ptr,
Scalar const * beta_ptr 
)
+ 		+inline
+
+ +
+
+

Member Function Documentation

+ +

◆ alpha()

+ +
+
+
+template<typename Scalar_ , typename FragmentMultiplyAdd_ = FragmentMultiplyAdd<Scalar_, Scalar_>>
+ + + + + +
+ + + + + + + +
CUTLASS_HOST_DEVICE Scalar cutlass::gemm::LinearScalingDevicePtr< Scalar_, FragmentMultiplyAdd_ >::Params::alpha () const
+ 		+inline
+
+ +
+
+ +

◆ beta()

+ +
+
+
+template<typename Scalar_ , typename FragmentMultiplyAdd_ = FragmentMultiplyAdd<Scalar_, Scalar_>>
+ + + + + +
+ + + + + + + +
CUTLASS_HOST_DEVICE Scalar cutlass::gemm::LinearScalingDevicePtr< Scalar_, FragmentMultiplyAdd_ >::Params::beta () const
+ 		+inline
+
+ +
+
+ +

◆ initialize() [1/3]

+ +
+
+
+template<typename Scalar_ , typename FragmentMultiplyAdd_ = FragmentMultiplyAdd<Scalar_, Scalar_>>
+ + + + + +
+ + + + + + + + + + + + + + + + + + +
CUTLASS_HOST_DEVICE int cutlass::gemm::LinearScalingDevicePtr< Scalar_, FragmentMultiplyAdd_ >::Params::initialize (Scalar alpha,
Scalar beta 
)
+ 		+inline
+
+ +
+
+ +

◆ initialize() [2/3]

+ +
+
+
+template<typename Scalar_ , typename FragmentMultiplyAdd_ = FragmentMultiplyAdd<Scalar_, Scalar_>>
+ + + + + +
+ + + + + + + + + + + + + + + + + + +
CUTLASS_HOST_DEVICE int cutlass::gemm::LinearScalingDevicePtr< Scalar_, FragmentMultiplyAdd_ >::Params::initialize (Scalar const * alpha,
Scalar const * beta 
)
+ 		+inline
+
+ +
+
+ +

◆ initialize() [3/3]

+ +
+
+
+template<typename Scalar_ , typename FragmentMultiplyAdd_ = FragmentMultiplyAdd<Scalar_, Scalar_>>
+
+template<typename GemmDesc_ >
+ + + + + +
+ + + + + + + + +
CUTLASS_HOST_DEVICE int cutlass::gemm::LinearScalingDevicePtr< Scalar_, FragmentMultiplyAdd_ >::Params::initialize (GemmDesc_ const & desc)
+ 		+inline
+
+ +
+
+
The documentation for this class was generated from the following file: +
+ +
+Generated on Wed Sep 19 2018 13:58:53 for Cutlass by   +doxygen + 1.8.14 +
+ + diff --git a/docs/classcutlass_1_1platform_1_1complex-members.html b/docs/classcutlass_1_1platform_1_1complex-members.html new file mode 100644 index 00000000..3e19742e --- /dev/null +++ b/docs/classcutlass_1_1platform_1_1complex-members.html @@ -0,0 +1,100 @@ + + + + + + + +Cutlass: Member List + + + + + + + + + + +
+
+ + + + + + +
+
Cutlass +
+
CUDA Templates for Linear Algebra Subroutines and Solvers
+
+
+ + + + + + + + +
+
+ + +
+ +
+ + +
+
+
+
cutlass::platform::complex< T > Member List
+
+
+ +

This is the complete list of members for cutlass::platform::complex< T >, including all inherited members.

+ + + + + + + + + + + +
complex(T r=T(0), T i=T(0))cutlass::platform::complex< T >inline
complex(cuFloatComplex const &z)cutlass::platform::complex< T >inline
complex(cuDoubleComplex const &z)cutlass::platform::complex< T >inline
imag() constcutlass::platform::complex< T >inline
imag()cutlass::platform::complex< T >inline
operator cuDoubleComplex() constcutlass::platform::complex< T >inline
operator cuFloatComplex() constcutlass::platform::complex< T >inline
real() constcutlass::platform::complex< T >inline
real()cutlass::platform::complex< T >inline
value_type typedefcutlass::platform::complex< T >
+ +
+Generated on Wed Sep 19 2018 13:58:53 for Cutlass by   +doxygen + 1.8.14 +
+ + diff --git a/docs/classcutlass_1_1platform_1_1complex.html b/docs/classcutlass_1_1platform_1_1complex.html new file mode 100644 index 00000000..672fef7e --- /dev/null +++ b/docs/classcutlass_1_1platform_1_1complex.html @@ -0,0 +1,413 @@ + + + + + + + +Cutlass: cutlass::platform::complex< T > Class Template Reference + + + + + + + + + + +
+
+ + + + + + +
+
Cutlass +
+
CUDA Templates for Linear Algebra Subroutines and Solvers
+
+
+ + + + + + + + +
+
+ + +
+ +
+ + +
+
+
+Public Types | +Public Member Functions | +List of all members
+
+
cutlass::platform::complex< T > Class Template Reference
+
+
+ +

#include <complex.h>

+ + + + + +

+Public Types

typedef T value_type
 Type alias for scalar type. More...
 
+ + + + + + + + + + + + + + + + + + + + + + + + + + + + +

+Public Member Functions

CUTLASS_HOST_DEVICE complex (T r=T(0), T i=T(0))
 Constructor. More...
 
CUTLASS_HOST_DEVICE complex (cuFloatComplex const &z)
 Conversion from cuFloatComplex. More...
 
CUTLASS_HOST_DEVICE complex (cuDoubleComplex const &z)
 Conversion from cuDoubleComplex. More...
 
CUTLASS_HOST_DEVICE T const & real () const
 Accesses the real part of the complex number. More...
 
CUTLASS_HOST_DEVICE T & real ()
 Accesses the real part of the complex number. More...
 
CUTLASS_HOST_DEVICE T const & imag () const
 Accesses the imaginary part of the complex number. More...
 
CUTLASS_HOST_DEVICE T & imag ()
 Accesses the imaginary part of the complex number. More...
 
CUTLASS_HOST_DEVICE operator cuFloatComplex () const
 Converts to cuFloatComplex. More...
 
CUTLASS_HOST_DEVICE operator cuDoubleComplex () const
 Converts to cuDoubleComplex. More...
 
+

Detailed Description

+

template<typename T>
+class cutlass::platform::complex< T >

+ +

Class for representing and manipulating complex numbers with conversions from built-in CUDA complex types.

+

Member Typedef Documentation

+ +

◆ value_type

+ +
+
+
+template<typename T>
+ + + + +
typedef T cutlass::platform::complex< T >::value_type
+
+ +
+
+

Constructor & Destructor Documentation

+ +

◆ complex() [1/3]

+ +
+
+
+template<typename T>
+ + + + + +
+ + + + + + + + + + + + + + + + + + +
CUTLASS_HOST_DEVICE cutlass::platform::complex< T >::complex (T r = T(0),
T i = T(0) 
)
+ 		+inline
+
+ +
+
+ +

◆ complex() [2/3]

+ +
+
+
+template<typename T>
+ + + + + +
+ + + + + + + + +
CUTLASS_HOST_DEVICE cutlass::platform::complex< T >::complex (cuFloatComplex const & z)
+ 		+inline
+
+ +
+
+ +

◆ complex() [3/3]

+ +
+
+
+template<typename T>
+ + + + + +
+ + + + + + + + +
CUTLASS_HOST_DEVICE cutlass::platform::complex< T >::complex (cuDoubleComplex const & z)
+ 		+inline
+
+ +
+
+

Member Function Documentation

+ +

◆ imag() [1/2]

+ +
+
+
+template<typename T>
+ + + + + +
+ + + + + + + +
CUTLASS_HOST_DEVICE T const& cutlass::platform::complex< T >::imag () const
+ 		+inline
+
+ +
+
+ +

◆ imag() [2/2]

+ +
+
+
+template<typename T>
+ + + + + +
+ + + + + + + +
CUTLASS_HOST_DEVICE T& cutlass::platform::complex< T >::imag ()
+ 		+inline
+
+ +
+
+ +

◆ operator cuDoubleComplex()

+ +
+
+
+template<typename T>
+ + + + + +
+ + + + + + + +
CUTLASS_HOST_DEVICE cutlass::platform::complex< T >::operator cuDoubleComplex () const
+ 		+inline
+
+ +
+
+ +

◆ operator cuFloatComplex()

+ +
+
+
+template<typename T>
+ + + + + +
+ + + + + + + +
CUTLASS_HOST_DEVICE cutlass::platform::complex< T >::operator cuFloatComplex () const
+ 		+inline
+
+ +
+
+ +

◆ real() [1/2]

+ +
+
+
+template<typename T>
+ + + + + +
+ + + + + + + +
CUTLASS_HOST_DEVICE T const& cutlass::platform::complex< T >::real () const
+ 		+inline
+
+ +
+
+ +

◆ real() [2/2]

+ +
+
+
+template<typename T>
+ + + + + +
+ + + + + + + +
CUTLASS_HOST_DEVICE T& cutlass::platform::complex< T >::real ()
+ 		+inline
+
+ +
+
+
The documentation for this class was generated from the following file: +
+ +
+Generated on Wed Sep 19 2018 13:58:53 for Cutlass by   +doxygen + 1.8.14 +
+ + diff --git a/docs/classcutlass_1_1platform_1_1unique__ptr-members.html b/docs/classcutlass_1_1platform_1_1unique__ptr-members.html index 696f4788..1242de68 100644 --- a/docs/classcutlass_1_1platform_1_1unique__ptr-members.html +++ b/docs/classcutlass_1_1platform_1_1unique__ptr-members.html @@ -98,7 +98,7 @@ $(function() {
-Generated on Thu May 3 2018 16:36:15 for Cutlass by   +Generated on Wed Sep 19 2018 13:58:53 for Cutlass by   doxygen 1.8.14
diff --git a/docs/classcutlass_1_1platform_1_1unique__ptr.html b/docs/classcutlass_1_1platform_1_1unique__ptr.html index cf455f2e..625e790b 100644 --- a/docs/classcutlass_1_1platform_1_1unique__ptr.html +++ b/docs/classcutlass_1_1platform_1_1unique__ptr.html @@ -546,7 +546,7 @@ template<class T, class Deleter = default_delete<T>>
-Generated on Thu May 3 2018 16:36:15 for Cutlass by   +Generated on Wed Sep 19 2018 13:58:53 for Cutlass by   doxygen 1.8.14
diff --git a/docs/classes.html b/docs/classes.html index 9896653f..6a517312 100644 --- a/docs/classes.html +++ b/docs/classes.html @@ -72,100 +72,116 @@ $(function() {
Class Index
-
a | b | c | d | e | f | g | h | i | l | m | n | p | r | s | t | u | v | w
+
a | b | c | d | e | f | g | h | i | k | l | m | n | p | r | s | t | u | v | w | z
- - - - - - - - - - - - - - - - + + + + + + + + + - - - + + + + + + + + + + - - + + + + - - - - - - - - - - - - - - - - - - - + + + + + + + + + + + + + + + + + + + + + + + + - - - + + + - - - - + + + + + + + - - - - - - + + + + + + + - - - - - - - + + + + + + + + + + + + +
  a  
-		FragmentMultiplyAdd (cutlass::gemm)   IgemmEpilogueScalar (cutlass::gemm)   Load< Scalar_, Lanes_, Memory_, true, 8 > (cutlass)   GlobalLoadStreamBase::SharedStorage (cutlass::gemm)   
FragmentMultiplyAdd< half > (cutlass::gemm)   IgemmEpilogueScalar< int > (cutlass::gemm)   log2_down (cutlass)   SimplifiedGemmEpilogueTraits (cutlass::gemm)   
aligned_chunk (cutlass::platform)   FragmentStore (cutlass)   IgemmEpilogueTraits (cutlass::gemm)   log2_down< N, 1, Count > (cutlass)   SimplifiedGemmTraits (cutlass::gemm)   
aligned_storage (cutlass::platform)   FragmentStore< IteratorFragment::kScalar, kAccessSize, Scalar_, Memory_, FragmentElement_, kStride > (cutlass)   IgemmEpilogueTraitsHelper (cutlass::gemm)   log2_up (cutlass)   SimplifiedGemmTraitsHelper (cutlass::gemm)   
AlignedStruct (cutlass)   FragmentStore< IteratorFragment::kWmmaMatrix, kAccessSize, Scalar_, Memory_, FragmentElement_, kStride > (cutlass)   IgemmFloatToInt8Converter (cutlass::gemm)   log2_up< N, 1, Count > (cutlass)   sqrt_est (cutlass)   
alignment_of (cutlass::platform)   
  g  
-		IgemmGlobalLoadTransformer (cutlass::gemm)   
  m  
-		StorageType (cutlass)   
alignment_of< const value_t > (cutlass::platform)   IgemmGlobalLoadTransformer< Fragment< int8_t, kElements_ >, float > (cutlass::gemm)   StorageType< 1 > (cutlass)   
alignment_of< const volatile value_t > (cutlass::platform)   Gemm (cutlass::gemm)   IgemmGlobalStoreTransformer (cutlass::gemm)   GemmTraits::MainLoopSharedStorage (cutlass::gemm)   StorageType< 2 > (cutlass)   
alignment_of< double2 > (cutlass::platform)   GemmConfig (cutlass::gemm)   IgemmGlobalStoreTransformer< float, Fragment< int8_t, kElements_ > > (cutlass::gemm)   MatrixLayout (cutlass)   StorageType< 4 > (cutlass)   
alignment_of< double4 > (cutlass::platform)   GemmDesc (cutlass::gemm)   IgemmInt8ToFloatConverter (cutlass::gemm)   MemorySpace (cutlass)   Store (cutlass)   
alignment_of< float4 > (cutlass::platform)   GemmEpilogue (cutlass::gemm)   IgemmSharedStoreTransformer (cutlass::gemm)   
  n  
-		Store< double, 2, Memory_, true, 16 > (cutlass)   
alignment_of< int4 > (cutlass::platform)   GemmEpilogueTraits (cutlass::gemm)   IgemmSwizzle (cutlass::gemm)   Store< Scalar_, Lanes_, Memory_, true, 16 > (cutlass)   
alignment_of< long4 > (cutlass::platform)   GemmEpilogueTraitsHelper (cutlass::gemm)   IgemmTileTraitsHelperA (cutlass::gemm)   nullptr_t (cutlass::platform)   Store< Scalar_, Lanes_, Memory_, true, 4 > (cutlass)   
alignment_of< longlong2 > (cutlass::platform)   GemmGlobalIteratorAb (cutlass::gemm)   IgemmTileTraitsHelperA< MatrixLayout::kColumnMajor, GemmConfig_ > (cutlass::gemm)   
  p  
-		Store< Scalar_, Lanes_, Memory_, true, 8 > (cutlass)   
alignment_of< longlong4 > (cutlass::platform)   GemmGlobalIteratorCd (cutlass::gemm)   IgemmTileTraitsHelperB (cutlass::gemm)   GemmTraits::StreamSharedStorage (cutlass::gemm)   
alignment_of< uint4 > (cutlass::platform)   GemmGlobalTileCdTraits (cutlass::gemm)   IgemmTileTraitsHelperB< MatrixLayout::kRowMajor, GemmConfig_ > (cutlass::gemm)   alignment_of::pad (cutlass::platform)   GemmEpilogueTraits::StreamSharedStorage (cutlass::gemm)   
alignment_of< ulong4 > (cutlass::platform)   GemmGlobalTileTraits (cutlass::gemm)   IgemmTraits (cutlass::gemm)   WmmaGemmGlobalIteratorCd::Params (cutlass::gemm)   
  t  
+		GemmConfig (cutlass::gemm)   IgemmTraitsHelper (cutlass::gemm)   LinearScalingDevicePtr::Params (cutlass::gemm)   Store< double, 2, Memory_, FragmentElementType::kScalar, double, kStride, 16 > (cutlass)   
GemmCoord (cutlass::gemm)   IgemmTransformerA (cutlass::gemm)   GlobalLoadStream::Params (cutlass::gemm)   Store< Scalar_, kAccessSize, Memory_, FragmentElementType::kScalar, Scalar_, 1, 2 > (cutlass)   
aligned_chunk (cutlass::platform)   GemmDesc (cutlass::gemm)   IgemmTransformerA< MatrixLayout::kColumnMajor, Iterator_ > (cutlass::gemm)   SharedStreamPair::Params (cutlass::gemm)   Store< Scalar_, kAccessSize, Memory_, FragmentElementType::kScalar, Scalar_, kStride, 16 > (cutlass)   
aligned_storage (cutlass::platform)   GemmEpilogue (cutlass::gemm)   IgemmTransformerA< MatrixLayout::kRowMajor, Iterator_ > (cutlass::gemm)   WmmaGemmGlobalIteratorCd::Params (cutlass::gemm)   Store< Scalar_, kAccessSize, Memory_, FragmentElementType::kScalar, Scalar_, kStride, 4 > (cutlass)   
AlignedStruct (cutlass)   GemmEpilogueTraits (cutlass::gemm)   IgemmTransformerB (cutlass::gemm)   ZipTileIterator::Params (cutlass)   Store< Scalar_, kAccessSize, Memory_, FragmentElementType::kScalar, Scalar_, kStride, 8 > (cutlass)   
alignment_of (cutlass::platform)   GemmEpilogueTraitsHelper (cutlass::gemm)   IgemmTransformerB< MatrixLayout::kColumnMajor, Iterator_ > (cutlass::gemm)   GemmTraits::Params (cutlass::gemm)   Store< Scalar_, kAccessSize, Memory_, FragmentElementType::kWmmaMatrix, FragmentElement_, kStride, size > (cutlass)   
alignment_of< const value_t > (cutlass::platform)   GemmGlobalIteratorAb (cutlass::gemm)   IgemmTransformerB< MatrixLayout::kRowMajor, Iterator_ > (cutlass::gemm)   LinearScaling::Params (cutlass::gemm)   GemmEpilogueTraits::StreamSharedStorage (cutlass::gemm)   
alignment_of< const volatile value_t > (cutlass::platform)   GemmGlobalIteratorCd (cutlass::gemm)   int4_t (cutlass)   GemmGlobalIteratorAb::Params (cutlass::gemm)   TensorRef< Storage_, Rank_, MapFunc_, 1, Index_, LongIndex_ >::StrideVector (cutlass)   
alignment_of< double2 > (cutlass::platform)   GemmGlobalTileCdTraits (cutlass::gemm)   integral_constant (cutlass::platform)   GlobalLoadStreamPair::Params (cutlass::gemm)   swizzleDirection (cutlass::gemm)   
alignment_of< double4 > (cutlass::platform)   GemmGlobalTileTraits (cutlass::gemm)   is_arithmetic (cutlass::platform)   GemmGlobalIteratorCd::Params (cutlass::gemm)   
  t  
alignment_of< ulonglong2 > (cutlass::platform)   GemmMultiplicandTraits (cutlass::gemm)   IgemmTraitsHelper (cutlass::gemm)   GemmTraits::Params (cutlass::gemm)   
alignment_of< ulonglong4 > (cutlass::platform)   GemmOperand (cutlass)   IgemmTransformerA (cutlass::gemm)   GlobalLoadStreamBase::Params (cutlass::gemm)   TensorRef (cutlass)   
alignment_of< volatile value_t > (cutlass::platform)   GemmOperandTraitsAb (cutlass::gemm)   IgemmTransformerA< MatrixLayout::kColumnMajor, Iterator_ > (cutlass::gemm)   TileIteratorBase::Params (cutlass)   TensorView (cutlass)   
alignment_of< float4 > (cutlass::platform)   GemmMultiplicandTraits (cutlass::gemm)   is_base_of (cutlass::platform)   GemmEpilogueTraits::Params (cutlass::gemm)   
alignment_of< int4 > (cutlass::platform)   GemmOperand (cutlass)   is_base_of_helper (cutlass::platform)   TileIteratorBase::Params (cutlass)   TensorRef (cutlass)   
alignment_of< long4 > (cutlass::platform)   GemmOperandTraitsAb (cutlass::gemm)   is_floating_point (cutlass::platform)   TileLoadIterator::Params (cutlass)   TensorRef< Storage_, Rank_, MapFunc_, 1, Index_, LongIndex_ > (cutlass)   
alignment_of< longlong2 > (cutlass::platform)   GemmSharedLoadTileATraits (cutlass::gemm)   is_fundamental (cutlass::platform)   TileStoreIterator::Params (cutlass)   TensorRefArray (cutlass)   
alignment_of< longlong4 > (cutlass::platform)   GemmSharedLoadTileBTraits (cutlass::gemm)   is_integral (cutlass::platform)   TileLoadStream::Params (cutlass)   TensorRefBatchStrided (cutlass)   
alignment_of< uint4 > (cutlass::platform)   GemmSharedLoadTileDTraits (cutlass::gemm)   is_integral< char > (cutlass::platform)   TileStoreStream::Params (cutlass)   TensorView (cutlass)   
alignment_of< ulong4 > (cutlass::platform)   GemmSharedStoreTileAbTraits (cutlass::gemm)   is_integral< const T > (cutlass::platform)   SharedLoadStream::Params (cutlass::gemm)   ThreadMultiplyAdd (cutlass::gemm)   
alignment_of< ulonglong2 > (cutlass::platform)   GemmSharedStoreTileDTraits (cutlass::gemm)   is_integral< const volatile T > (cutlass::platform)   plus (cutlass::platform)   ThreadMultiplyAdd< ThreadGemmShape_, ThreadsPerWarp_, half, half, float > (cutlass::gemm)   
alignment_of< ulonglong4 > (cutlass::platform)   GemmSharedStoreWithSkewTileAbTraits (cutlass::gemm)   is_integral< int > (cutlass::platform)   PredicatedTileLoadStream (cutlass)   ThreadMultiplyAdd< ThreadGemmShape_, ThreadsPerWarp_, half, half, half > (cutlass::gemm)   
alignment_of< volatile value_t > (cutlass::platform)   GemmTileTraitsHelperA (cutlass::gemm)   is_integral< long > (cutlass::platform)   PredicatedTileStoreStream (cutlass)   ThreadMultiplyAdd< ThreadGemmShape_, ThreadsPerWarp_, int8_t, int8_t, int > (cutlass::gemm)   
  b  
-		GemmSharedLoadTileATraits (cutlass::gemm)   IgemmTransformerA< MatrixLayout::kRowMajor, Iterator_ > (cutlass::gemm)   GemmGlobalIteratorCd::Params (cutlass::gemm)   ThreadMultiplyAdd (cutlass::gemm)   
GemmSharedLoadTileBTraits (cutlass::gemm)   IgemmTransformerB (cutlass::gemm)   TileLoadIterator::Params (cutlass)   ThreadMultiplyAdd< AccumulatorsPerThread_, ThreadsPerWarp_, half, half, half > (cutlass::gemm)   
bool_constant (cutlass::platform)   GemmSharedLoadTileDTraits (cutlass::gemm)   IgemmTransformerB< MatrixLayout::kColumnMajor, Iterator_ > (cutlass::gemm)   TileStoreIterator::Params (cutlass)   ThreadMultiplyAdd< AccumulatorsPerThread_, ThreadsPerWarp_, int8_t, int8_t, int > (cutlass::gemm)   
GemmTileTraitsHelperA< MatrixLayout::kColumnMajor, GemmConfig_ > (cutlass::gemm)   is_integral< long long > (cutlass::platform)   PredicateTileAdapter (cutlass)   GemmSharedStoreTileAbTraits::ThreadOffset (cutlass::gemm)   
GemmTileTraitsHelperA< MatrixLayout::kRowMajor, GemmConfig_ > (cutlass::gemm)   is_integral< short > (cutlass::platform)   TileLoadStream::PredicateVector (cutlass)   WmmaGemmGlobalIteratorCdTraits::ThreadOffset (cutlass::gemm)   
bin1_t (cutlass)   GemmTileTraitsHelperB (cutlass::gemm)   is_integral< signed char > (cutlass::platform)   PredicateVector (cutlass)   GemmGlobalTileCdTraits::ThreadOffset (cutlass::gemm)   
bool_constant (cutlass::platform)   GemmTileTraitsHelperB< MatrixLayout::kColumnMajor, GemmConfig_ > (cutlass::gemm)   is_integral< unsigned char > (cutlass::platform)   TileStoreStream::PredicateVector (cutlass)   GemmSharedLoadTileATraits::ThreadOffset (cutlass::gemm)   
  c  
-		GemmSharedStoreTileAbTraits (cutlass::gemm)   IgemmTransformerB< MatrixLayout::kRowMajor, Iterator_ > (cutlass::gemm)   GemmEpilogueTraits::Params (cutlass::gemm)   GemmSharedLoadTileBTraits::ThreadOffset (cutlass::gemm)   
GemmSharedStoreTileDTraits (cutlass::gemm)   integral_constant (cutlass::platform)   Gemm::Params (cutlass::gemm)   GemmGlobalTileCdTraits::ThreadOffset (cutlass::gemm)   
ClearAccumulators (cutlass::gemm)   GemmSharedStoreWithSkewTileAbTraits (cutlass::gemm)   is_arithmetic (cutlass::platform)   SharedLoadStream::Params (cutlass::gemm)   IgemmContiguousGlobalTileTraits::ThreadOffset (cutlass::gemm)   
ComputeOffsetFromShape (cutlass)   GemmTileTraitsHelperA (cutlass::gemm)   is_base_of (cutlass::platform)   LinearScaling::Params (cutlass::gemm)   GemmGlobalTileTraits::ThreadOffset (cutlass::gemm)   
ComputeOffsetFromShape< Shape< 1, kSh_, kSw_, 1 > > (cutlass)   GemmTileTraitsHelperA< MatrixLayout::kColumnMajor, GemmConfig_ > (cutlass::gemm)   is_base_of_helper (cutlass::platform)   GemmGlobalIteratorAb::Params (cutlass::gemm)   GemmSharedLoadTileDTraits::ThreadOffset (cutlass::gemm)   
ComputeOffsetFromShape< Shape< 1, kSh_, kSw_, kSc_ > > (cutlass)   GemmTileTraitsHelperA< MatrixLayout::kRowMajor, GemmConfig_ > (cutlass::gemm)   is_floating_point (cutlass::platform)   plus (cutlass::platform)   GemmSharedLoadTileATraits::ThreadOffset (cutlass::gemm)   
ComputeOffsetFromStrides (cutlass)   GemmTileTraitsHelperB (cutlass::gemm)   is_fundamental (cutlass::platform)   PredicateTileAdapter (cutlass)   GemmSharedStoreTileDTraits::ThreadOffset (cutlass::gemm)   
ComputeOffsetFromStrides< Shape< 1, S_h_, S_w_, 1 > > (cutlass)   GemmTileTraitsHelperB< MatrixLayout::kColumnMajor, GemmConfig_ > (cutlass::gemm)   is_integral (cutlass::platform)   PredicateVector (cutlass)   HgemmCrosswiseGlobalTileTraits::ThreadOffset (cutlass::gemm)   
ComputeOffsetFromStrides< Shape< 1, S_h_, S_w_, S_c_ > > (cutlass)   GemmTileTraitsHelperB< MatrixLayout::kRowMajor, GemmConfig_ > (cutlass::gemm)   is_integral< char > (cutlass::platform)   ProjectOperand (cutlass::gemm)   GemmSharedStoreTileAbTraits::ThreadOffset (cutlass::gemm)   
ComputeThreadOffsetFromStrides (cutlass)   GemmTraits (cutlass::gemm)   is_integral< const T > (cutlass::platform)   ProjectOperand< GemmOperand::kA, Kstrided > (cutlass::gemm)   TileTraitsWarpRake::ThreadOffset (cutlass)   
ComputeThreadOffsetFromStrides< Shape< 1, T_h_, T_w_, 1 >, Shape< 1, S_h_, S_w_, 1 > > (cutlass)   GetExtent (cutlass::gemm)   is_integral< const volatile T > (cutlass::platform)   ProjectOperand< GemmOperand::kB, Kstrided > (cutlass::gemm)   GemmSharedStoreWithSkewTileAbTraits::ThreadOffset (cutlass::gemm)   
ComputeThreadOffsetFromStrides< Shape< 1, T_h_, T_w_, T_c_ >, Shape< 1, S_h_, S_w_, S_c_ > > (cutlass)   GetExtent< GemmOperand::kA, Tile_ > (cutlass::gemm)   is_integral< int > (cutlass::platform)   ProjectOperand< GemmOperand::kC, true > (cutlass::gemm)   WmmaGemmGlobalIteratorCdTraits::ThreadOffset (cutlass::gemm)   
conditional (cutlass::platform)   GetExtent< GemmOperand::kB, Tile_ > (cutlass::gemm)   is_integral< long > (cutlass::platform)   ProjectOperand< GemmOperand::kD, true > (cutlass::gemm)   TiledThreadOffset (cutlass)   
conditional< false, T, F > (cutlass::platform)   GemmTraits::GlobalLoadStream (cutlass::gemm)   is_integral< long long > (cutlass::platform)   
  r  
-		TileIteratorBase (cutlass)   
PredicateVector::ConstIterator (cutlass)   GlobalLoadStream (cutlass::gemm)   is_integral< short > (cutlass::platform)   TileLoadIterator (cutlass)   
ConstPredicateTileAdapter (cutlass)   GlobalLoadStreamBase (cutlass::gemm)   is_integral< signed char > (cutlass::platform)   remove_const (cutlass::platform)   TileStoreIterator (cutlass)   
Convert (cutlass)   greater (cutlass::platform)   is_integral< unsigned char > (cutlass::platform)   remove_const< const T > (cutlass::platform)   TileTraits (cutlass)   
Convert< Fragment< InputScalar_, kScalars_ >, Fragment< OutputScalar_, kScalars_ > > (cutlass)   
  h  
-		is_integral< unsigned int > (cutlass::platform)   remove_cv (cutlass::platform)   TileTraitsContiguousMajor (cutlass)   
Coord (cutlass)   is_integral< unsigned long > (cutlass::platform)   remove_volatile (cutlass::platform)   TileTraitsStandard (cutlass)   
Copy (cutlass)   HgemmConfig (cutlass::gemm)   is_integral< unsigned long long > (cutlass::platform)   remove_volatile< volatile T > (cutlass::platform)   TileTraitsStrideMajor (cutlass)   
GemmTileTraitsHelperB< MatrixLayout::kRowMajor, GemmConfig_ > (cutlass::gemm)   is_integral< unsigned int > (cutlass::platform)   ProjectOperand (cutlass::gemm)   GemmSharedStoreWithSkewTileAbTraits::ThreadOffset (cutlass::gemm)   
GemmTraits (cutlass::gemm)   is_integral< unsigned long > (cutlass::platform)   ProjectOperand< GemmOperand::kA, Kstrided > (cutlass::gemm)   IgemmGlobalTileTraits::ThreadOffset (cutlass::gemm)   
ClearAccumulators (cutlass::gemm)   GetExtent (cutlass::gemm)   is_integral< unsigned long long > (cutlass::platform)   ProjectOperand< GemmOperand::kB, Kstrided > (cutlass::gemm)   GemmSharedLoadTileBTraits::ThreadOffset (cutlass::gemm)   
MatrixLayout::ColumnMajor (cutlass)   GetExtent< GemmOperand::kA, Tile_ > (cutlass::gemm)   is_integral< unsigned short > (cutlass::platform)   ProjectOperand< GemmOperand::kC, true > (cutlass::gemm)   GemmGlobalTileTraits::ThreadOffset (cutlass::gemm)   
MatrixLayout::ColumnMajorBlockLinear (cutlass)   GetExtent< GemmOperand::kB, Tile_ > (cutlass::gemm)   is_integral< volatile T > (cutlass::platform)   ProjectOperand< GemmOperand::kD, true > (cutlass::gemm)   GemmSharedLoadTileDTraits::ThreadOffset (cutlass::gemm)   
ColumnMajorBlockSwizzle (cutlass::gemm)   GlobalLoadStream (cutlass::gemm)   is_pointer (cutlass::platform)   
  r  
+		TileTraitsWarpRake::ThreadOffset (cutlass)   
MatrixLayout::ColumnMajorInterleaved (cutlass)   GlobalLoadStreamPair (cutlass::gemm)   is_pointer_helper (cutlass::platform)   GemmSharedStoreTileDTraits::ThreadOffset (cutlass::gemm)   
complex (cutlass::platform)   greater (cutlass::platform)   is_pointer_helper< T * > (cutlass::platform)   RegularTilePredicateFunctor (cutlass)   HgemmCrosswiseGlobalTileTraits::ThreadOffset (cutlass::gemm)   
ComputeOffsetFromShape (cutlass)   
  h  
+		is_pow2 (cutlass)   remove_const (cutlass::platform)   TileAllocation (cutlass)   
ComputeOffsetFromStrides (cutlass)   is_same (cutlass::platform)   remove_const< const T > (cutlass::platform)   TileCoord (cutlass)   
ComputeThreadOffsetFromStrides (cutlass)   HgemmConfig (cutlass::gemm)   is_same< A, A > (cutlass::platform)   remove_cv (cutlass::platform)   TiledThreadOffset (cutlass)   
ComputeThreadOffsetFromStrides< Shape< 1, T_h_, T_w_, 1 >, Shape< 1, S_h_, S_w_, 1 > > (cutlass)   HgemmCrosswiseGlobalTileTraits (cutlass::gemm)   is_trivially_copyable (cutlass::platform)   remove_volatile (cutlass::platform)   TileIteratorBase (cutlass)   
ComputeThreadOffsetFromStrides< Shape< 1, T_h_, T_w_, T_c_ >, Shape< 1, S_h_, S_w_, S_c_ > > (cutlass)   HgemmSwizzle (cutlass::gemm)   is_void (cutlass::platform)   remove_volatile< volatile T > (cutlass::platform)   TileLoadIterator (cutlass)   
conditional (cutlass::platform)   HgemmTileTraitsHelperA (cutlass::gemm)   is_volatile (cutlass::platform)   ReshapeThreads (cutlass::gemm)   TileLoadStream (cutlass)   
conditional< false, T, F > (cutlass::platform)   HgemmTileTraitsHelperA< MatrixLayout::kRowMajor, GemmConfig_ > (cutlass::gemm)   is_volatile< volatile T > (cutlass::platform)   ReshapeThreads< Tile_, Threads_, true > (cutlass::gemm)   TileStoreIterator (cutlass)   
PredicateVector::ConstIterator (cutlass)   HgemmTileTraitsHelperB (cutlass::gemm)   PredicateVector::Iterator (cutlass)   ReshapeTile (cutlass)   TileStoreStream (cutlass)   
TensorRefBatchStrided::ConstIterator (cutlass)   HgemmTileTraitsHelperB< MatrixLayout::kColumnMajor, GemmConfig_ > (cutlass::gemm)   IteratorAdvance (cutlass)   ReshapeTile< Tile_, kAccessSize_, true > (cutlass)   TileTraits (cutlass)   
TensorRefArray::ConstIterator (cutlass)   HgemmTraits (cutlass::gemm)   
  k  
+		MatrixLayout::RowMajor (cutlass)   TileTraitsContiguousMajor (cutlass)   
ConstPredicateTileAdapter (cutlass)   HgemmTraitsHelper (cutlass::gemm)   MatrixLayout::RowMajorBlockLinear (cutlass)   TileTraitsStandard (cutlass)   
MatrixLayout::ContiguousLayout (cutlass)   HgemmTransformerA (cutlass::gemm)   KernelLaunchConfiguration (cutlass)   RowMajorBlockSwizzle (cutlass::gemm)   TileTraitsStrideMajor (cutlass)   
Convert (cutlass)   HgemmTransformerA< MatrixLayout::kColumnMajor, Iterator_ > (cutlass::gemm)   
  l  
+		MatrixLayout::RowMajorInterleaved (cutlass)   TileTraitsWarpRake (cutlass)   
Convert< Fragment< InputScalar_, kScalars_ >, Fragment< OutputScalar_, kScalars_ > > (cutlass)   HgemmTransformerA< MatrixLayout::kRowMajor, Iterator_ > (cutlass::gemm)   
  s  
+		PredicateVector::TrivialIterator (cutlass)   
Coord (cutlass)   HgemmTransformerB (cutlass::gemm)   Launch (cutlass::gemm)   TrivialPredicateTileAdapter (cutlass)   
Copy (cutlass)   HgemmTransformerB< MatrixLayout::kColumnMajor, Iterator_ > (cutlass::gemm)   Launch< Gemm, false > (cutlass::gemm)   ScalarIO (cutlass)   
  u  
+
  d  
-		HgemmCrosswiseGlobalTileTraits (cutlass::gemm)   is_integral< unsigned short > (cutlass::platform)   ReshapeThreads (cutlass::gemm)   TileTraitsWarpRake (cutlass)   
HgemmSwizzle (cutlass::gemm)   is_integral< volatile T > (cutlass::platform)   ReshapeThreads< Tile_, Threads_, true > (cutlass::gemm)   PredicateVector::TrivialIterator (cutlass)   
default_delete (cutlass::platform)   HgemmTileTraitsHelperA (cutlass::gemm)   is_pointer (cutlass::platform)   ReshapeTile (cutlass)   TrivialPredicateTileAdapter (cutlass)   
default_delete< T[]> (cutlass::platform)   HgemmTileTraitsHelperA< MatrixLayout::kRowMajor, GemmConfig_ > (cutlass::gemm)   is_pointer_helper (cutlass::platform)   ReshapeTile< Tile_, kAccessSize_, true > (cutlass)   
  u  
+		HgemmTransformerB< MatrixLayout::kRowMajor, Iterator_ > (cutlass::gemm)   less (cutlass::platform)   ScalarOrPointer (cutlass::detail)   
  i  
+		LinearScaling (cutlass::gemm)   SgemmConfig (cutlass::gemm)   uint4_t (cutlass)   
DebugType   LinearScalingDevicePtr (cutlass::gemm)   SgemmLBTraits (cutlass::gemm)   unique_ptr (cutlass::platform)   
DebugValue   Identity (cutlass)   Load (cutlass)   SgemmTraits (cutlass::gemm)   
  v  
DgemmConfig (cutlass::gemm)   HgemmTileTraitsHelperB (cutlass::gemm)   is_pointer_helper< T * > (cutlass::platform)   
  s  
-
DgemmTraits (cutlass::gemm)   HgemmTileTraitsHelperB< MatrixLayout::kColumnMajor, GemmConfig_ > (cutlass::gemm)   is_pow2 (cutlass)   unique_ptr (cutlass::platform)   
divide_assert (cutlass)   HgemmTraits (cutlass::gemm)   is_same (cutlass::platform)   SgemmConfig (cutlass::gemm)   
  v  
-
is_base_of_helper::dummy (cutlass::platform)   HgemmTraitsHelper (cutlass::gemm)   is_same< A, A > (cutlass::platform)   SgemmTraits (cutlass::gemm)   
default_delete (cutlass::platform)   IdentityBlockSwizzle (cutlass::gemm)   Load< double, 2, Memory_, FragmentElementType::kScalar, double, kStride, 16 > (cutlass)   Shape (cutlass)   
default_delete< T[]> (cutlass::platform)   IdentityTensorMapFunc (cutlass)   Load< Scalar_, kAccessSize, Memory_, FragmentElementType::kScalar, Scalar_, 1, 2 > (cutlass)   ShapeAdd (cutlass)   Vector (cutlass)   
DgemmConfig (cutlass::gemm)   IgemmConfig (cutlass::gemm)   Load< Scalar_, kAccessSize, Memory_, FragmentElementType::kScalar, Scalar_, kStride, 16 > (cutlass)   ShapeCount (cutlass)   Vector< bin1_t, kLanes_ > (cutlass)   
DgemmTraits (cutlass::gemm)   IgemmConfig< OutputTile_, int8_t, ThreadGemmShape_ > (cutlass::gemm)   Load< Scalar_, kAccessSize, Memory_, FragmentElementType::kScalar, Scalar_, kStride, 4 > (cutlass)   ShapeDiv (cutlass)   Vector< half, 1 > (cutlass)   
divide_assert (cutlass)   IgemmEpilogue (cutlass::gemm)   Load< Scalar_, kAccessSize, Memory_, FragmentElementType::kScalar, Scalar_, kStride, 8 > (cutlass)   ShapeDivCeiling (cutlass)   Vector< half, kLanes_ > (cutlass)   
is_base_of_helper::dummy (cutlass::platform)   IgemmEpilogue< GemmEpilogueTraits_, true > (cutlass::gemm)   Load< Scalar_, kAccessSize, Memory_, FragmentElementType::kWmmaMatrix, FragmentElement_, kStride, size > (cutlass)   ShapeMax (cutlass)   Vector< int4_t, kLanes_ > (cutlass)   
DumpType (cutlass)   IgemmEpilogueScalar (cutlass::gemm)   Load< Vector< bin1_t, 32 >, kAccessSize, Memory_, FragmentElementType::kWmmaMatrix, FragmentElement_, kStride, size > (cutlass)   ShapeMin (cutlass)   Vector< uint4_t, kLanes_ > (cutlass)   
  e  
-		HgemmTransformerA (cutlass::gemm)   is_trivially_copyable (cutlass::platform)   Shape (cutlass)   Vector (cutlass)   
HgemmTransformerA< MatrixLayout::kColumnMajor, Iterator_ > (cutlass::gemm)   is_void (cutlass::platform)   ShapeAdd (cutlass)   Vector< half, kLanes_ > (cutlass)   
enable_if (cutlass::platform)   HgemmTransformerA< MatrixLayout::kRowMajor, Iterator_ > (cutlass::gemm)   is_volatile (cutlass::platform)   ShapeCount (cutlass)   Vectorize (cutlass)   
enable_if< false, T > (cutlass::platform)   HgemmTransformerB (cutlass::gemm)   is_volatile< volatile T > (cutlass::platform)   ShapeDiv (cutlass)   Vectorize< Element_, 1 > (cutlass)   
Extent (cutlass)   HgemmTransformerB< MatrixLayout::kColumnMajor, Iterator_ > (cutlass::gemm)   PredicateVector::Iterator (cutlass)   ShapeMax (cutlass)   VectorTraits (cutlass)   
Extent< Vector< T, Lanes > > (cutlass)   HgemmTransformerB< MatrixLayout::kRowMajor, Iterator_ > (cutlass::gemm)   IteratorAdvance (cutlass)   ShapeMin (cutlass)   VectorTraits< Vector< T, Lanes > > (cutlass)   
Extent< Vector< T, Lanes > const > (cutlass)   
  i  
-		IteratorFragment (cutlass)   ShapeMul (cutlass)   VectorTraits< Vector< T, Lanes > const > (cutlass)   
IgemmEpilogueScalar< int > (cutlass::gemm)   Load< Vector< int4_t, 8 >, kAccessSize, Memory_, FragmentElementType::kWmmaMatrix, FragmentElement_, kStride, size > (cutlass)   ShapeMul (cutlass)   Vectorize (cutlass)   
IgemmEpilogueTraits (cutlass::gemm)   Load< Vector< uint4_t, 8 >, kAccessSize, Memory_, FragmentElementType::kWmmaMatrix, FragmentElement_, kStride, size > (cutlass)   ShapeScale (cutlass)   Vectorize< Vector< bin1_t, 32 >, kLanes_ > (cutlass)   
enable_if (cutlass::platform)   IgemmEpilogueTraitsHelper (cutlass::gemm)   log2_down (cutlass)   ShapeStrides (cutlass)   Vectorize< Vector< int4_t, 8 >, kLanes_ > (cutlass)   
enable_if< false, T > (cutlass::platform)   IgemmFloatToInt8Converter (cutlass::gemm)   log2_down< N, 1, Count > (cutlass)   ShapeSub (cutlass)   Vectorize< Vector< uint4_t, 8 >, kLanes_ > (cutlass)   
Extent (cutlass)   IgemmGlobalIteratorAb (cutlass::gemm)   log2_up (cutlass)   SharedLoadStream (cutlass::gemm)   VectorTraits (cutlass)   
Extent< Vector< T, Lanes > > (cutlass)   IgemmGlobalLoadTransformer (cutlass::gemm)   log2_up< N, 1, Count > (cutlass)   GemmEpilogueTraits::SharedStorage (cutlass::gemm)   VectorTraits< Vector< T, Lanes > > (cutlass)   
Extent< Vector< T, Lanes > const > (cutlass)   IgemmGlobalLoadTransformer< Fragment< int8_t, kElements_ >, float > (cutlass::gemm)   
  m  
+		GlobalLoadStreamPair::SharedStorage (cutlass::gemm)   VectorTraits< Vector< T, Lanes > const > (cutlass)   
  f  
-
  l  
-		ShapeScale (cutlass)   
  w  
+		IgemmGlobalStoreTransformer (cutlass::gemm)   GemmTraits::SharedStorage (cutlass::gemm)   
  w  
Identity (cutlass)   ShapeStrides (cutlass)   
Fragment (cutlass)   IdentityBlockSwizzle (cutlass::gemm)   less (cutlass::platform)   ShapeSub (cutlass)   WmmaGemmGlobalIteratorCd (cutlass::gemm)   
FragmentConstIterator (cutlass)   IgemmConfig (cutlass::gemm)   LinearScaling (cutlass::gemm)   GemmTraits::SharedLoadStream (cutlass::gemm)   WmmaGemmGlobalIteratorCdTraits (cutlass::gemm)   
FragmentIterator (cutlass)   IgemmConfig< OutputTile_, int8_t, AccumulatorsPerThread_ > (cutlass::gemm)   Load (cutlass)   SharedLoadStream (cutlass::gemm)   
FragmentLoad (cutlass)   IgemmContiguousGlobalTileTraits (cutlass::gemm)   Load< double, 2, Memory_, true, 16 > (cutlass)   ClearAccumulators::SharedStorage (cutlass::gemm)   
FragmentLoad< IteratorFragment::kScalar, kAccessSize, Scalar_, Memory_, FragmentElement_, kStride > (cutlass)   IgemmEpilogue (cutlass::gemm)   Load< Scalar_, Lanes_, Memory_, true, 16 > (cutlass)   GemmEpilogueTraits::SharedStorage (cutlass::gemm)   
FragmentLoad< IteratorFragment::kWmmaMatrix, kAccessSize, Scalar_, Memory_, FragmentElement_, kStride > (cutlass)   IgemmEpilogue< GemmEpilogueTraits_, true > (cutlass::gemm)   Load< Scalar_, Lanes_, Memory_, true, 4 > (cutlass)   GemmTraits::SharedStorage (cutlass::gemm)   
IgemmGlobalStoreTransformer< float, Fragment< int8_t, kElements_ > > (cutlass::gemm)   GemmTraits::MainLoopSharedStorage (cutlass::gemm)   GlobalLoadStream::SharedStorage (cutlass::gemm)   
Fp16SgemmConfig (cutlass::gemm)   IgemmGlobalTileTraits (cutlass::gemm)   MatrixCoord (cutlass)   ClearAccumulators::SharedStorage (cutlass::gemm)   WmmaGemmGlobalIteratorCd (cutlass::gemm)   
Fp16SgemmSgemmTraits (cutlass::gemm)   IgemmInt8ToFloatConverter (cutlass::gemm)   MatrixLayout (cutlass)   SharedStreamPair (cutlass::gemm)   WmmaGemmGlobalIteratorCdTraits (cutlass::gemm)   
Fragment (cutlass)   IgemmSharedStoreTransformer (cutlass::gemm)   MatrixTransform (cutlass)   SimplifiedGemmEpilogueTraits (cutlass::gemm)   
  z  
+
FragmentConstIterator (cutlass)   IgemmSwizzle (cutlass::gemm)   Max (cutlass)   SimplifiedGemmTraits (cutlass::gemm)   
FragmentElementType (cutlass)   IgemmTileTraitsHelperA (cutlass::gemm)   MemorySpace (cutlass)   SimplifiedGemmTraitsHelper (cutlass::gemm)   ZipConvert (cutlass)   
FragmentIterator (cutlass)   IgemmTileTraitsHelperA< MatrixLayout::kColumnMajor, GemmConfig_, Index_ > (cutlass::gemm)   Min (cutlass)   sqrt_est (cutlass)   ZipFragment (cutlass)   
FragmentMultiplyAdd (cutlass::gemm)   IgemmTileTraitsHelperA< MatrixLayout::kRowMajor, GemmConfig_, Index_ > (cutlass::gemm)   
  n  
+		StorageType (cutlass)   ZipTensorRef (cutlass)   
FragmentMultiplyAdd< half, half, true > (cutlass::gemm)   IgemmTileTraitsHelperB (cutlass::gemm)   StorageType< 1 > (cutlass)   ZipTileAllocation (cutlass)   
  g  
+		IgemmTileTraitsHelperB< MatrixLayout::kColumnMajor, GemmConfig_, Index_ > (cutlass::gemm)   nullptr_t (cutlass::platform)   StorageType< 2 > (cutlass)   ZipTileIterator (cutlass)   
IgemmTileTraitsHelperB< MatrixLayout::kRowMajor, GemmConfig_, Index_ > (cutlass::gemm)   
  p  
+		StorageType< 4 > (cutlass)   
Gemm (cutlass::gemm)   IgemmTraits (cutlass::gemm)   Store (cutlass)   
alignment_of::pad (cutlass::platform)   
-
a | b | c | d | e | f | g | h | i | l | m | n | p | r | s | t | u | v | w
+
a | b | c | d | e | f | g | h | i | k | l | m | n | p | r | s | t | u | v | w | z
-Generated on Thu May 3 2018 16:36:15 for Cutlass by   +Generated on Wed Sep 19 2018 13:58:53 for Cutlass by   doxygen 1.8.14
diff --git a/docs/clear__accumulators_8h.html b/docs/clear__accumulators_8h.html index b4bd3b39..cd8f6307 100644 --- a/docs/clear__accumulators_8h.html +++ b/docs/clear__accumulators_8h.html @@ -82,7 +82,7 @@ $(function() {

Defines abstractions for efficiently clearing accumulator tiles. More...

-
#include <cutlass/vector.h>
+
#include "cutlass/vector.h"

Go to the source code of this file.

@@ -104,7 +104,7 @@ Namespaces
-Generated on Thu May 3 2018 16:36:12 for Cutlass by   +Generated on Wed Sep 19 2018 13:58:50 for Cutlass by   doxygen 1.8.14
diff --git a/docs/clear__accumulators_8h_source.html b/docs/clear__accumulators_8h_source.html index 1a6f517f..7c0423a5 100644 --- a/docs/clear__accumulators_8h_source.html +++ b/docs/clear__accumulators_8h_source.html @@ -76,16 +76,17 @@ $(function() {
clear_accumulators.h
-Go to the documentation of this file.
1 /***************************************************************************************************
2  * Copyright (c) 2017-2018, NVIDIA CORPORATION. All rights reserved.
3  *
4  * Redistribution and use in source and binary forms, with or without modification, are permitted
5  * provided that the following conditions are met:
6  * * Redistributions of source code must retain the above copyright notice, this list of
7  * conditions and the following disclaimer.
8  * * Redistributions in binary form must reproduce the above copyright notice, this list of
9  * conditions and the following disclaimer in the documentation and/or other materials
10  * provided with the distribution.
11  * * Neither the name of the NVIDIA CORPORATION nor the names of its contributors may be used
12  * to endorse or promote products derived from this software without specific prior written
13  * permission.
14  *
15  * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND ANY EXPRESS OR
16  * IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND
17  * FITNESS FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE
18  * FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING,
19  * BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS;
20  * OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT,
21  * STRICT LIABILITY, OR TOR (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
22  * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
23  *
24  **************************************************************************************************/
28 #pragma once
29 
30 #include <cutlass/vector.h>
31 
32 namespace cutlass {
33 namespace gemm {
34 
36 
37 template <typename Scalar_, int kLanes_ = 1>
38 struct ClearAccumulators {
40  struct SharedStorage {};
41 
43  CUTLASS_DEVICE ClearAccumulators(SharedStorage& shared_storage) {}
44 
46  template <typename Fragment_>
47  CUTLASS_DEVICE void clear(Fragment_& fragment) {
48  fragment.clear();
49  }
50 };
51 
53 
54 } // namespace gemm
55 } // namespace cutlass
cutlass
Definition: convert.h:33
+Go to the documentation of this file.
1 /***************************************************************************************************
2  * Copyright (c) 2017-2018, NVIDIA CORPORATION. All rights reserved.
3  *
4  * Redistribution and use in source and binary forms, with or without modification, are permitted
5  * provided that the following conditions are met:
6  * * Redistributions of source code must retain the above copyright notice, this list of
7  * conditions and the following disclaimer.
8  * * Redistributions in binary form must reproduce the above copyright notice, this list of
9  * conditions and the following disclaimer in the documentation and/or other materials
10  * provided with the distribution.
11  * * Neither the name of the NVIDIA CORPORATION nor the names of its contributors may be used
12  * to endorse or promote products derived from this software without specific prior written
13  * permission.
14  *
15  * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND ANY EXPRESS OR
16  * IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND
17  * FITNESS FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE
18  * FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING,
19  * BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS;
20  * OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT,
21  * STRICT LIABILITY, OR TOR (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
22  * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
23  *
24  **************************************************************************************************/
28 #pragma once
29 
30 #include "cutlass/vector.h"
31 
32 namespace cutlass {
33 namespace gemm {
34 
36 
37 template <typename Scalar_, int kLanes_ = 1>
38 struct ClearAccumulators {
40  struct SharedStorage {};
41 
43  CUTLASS_DEVICE ClearAccumulators(SharedStorage& shared_storage) {}
44 
46  CUTLASS_DEVICE ClearAccumulators() {}
47 
49  template <typename Fragment_>
50  CUTLASS_DEVICE void clear(Fragment_& fragment) {
51  fragment.clear();
52  }
53 };
54 
56 
57 } // namespace gemm
58 } // namespace cutlass
cutlass
Definition: convert.h:33
cutlass::gemm::ClearAccumulators
Definition: clear_accumulators.h:38
cutlass::gemm::ClearAccumulators::ClearAccumulators
CUTLASS_DEVICE ClearAccumulators(SharedStorage &shared_storage)
Ctor.
Definition: clear_accumulators.h:43
vector.h
Defines a 1D vector of elements held in the registers of each thread.
-
cutlass::gemm::ClearAccumulators::clear
CUTLASS_DEVICE void clear(Fragment_ &fragment)
Clear the fragment.
Definition: clear_accumulators.h:47
+
cutlass::gemm::ClearAccumulators::clear
CUTLASS_DEVICE void clear(Fragment_ &fragment)
Clear the fragment.
Definition: clear_accumulators.h:50
cutlass::gemm::ClearAccumulators::SharedStorage
The shared storage.
Definition: clear_accumulators.h:40
+
cutlass::gemm::ClearAccumulators::ClearAccumulators
CUTLASS_DEVICE ClearAccumulators()
Ctor.
Definition: clear_accumulators.h:46
-Generated on Thu May 3 2018 16:36:12 for Cutlass by   +Generated on Wed Sep 19 2018 13:58:50 for Cutlass by   doxygen 1.8.14
diff --git a/docs/complex_8h.html b/docs/complex_8h.html new file mode 100644 index 00000000..e94494d2 --- /dev/null +++ b/docs/complex_8h.html @@ -0,0 +1,263 @@ + + + + + + + +Cutlass: complex.h File Reference + + + + + + + + + + +
+
+
+ + + + + +
+
Cutlass +
+
CUDA Templates for Linear Algebra Subroutines and Solvers
+
+
+ + + + + + + + +
+
+ + +
+ +
+ + + +
+
+Classes | +Namespaces | +Functions
+
+
complex.h File Reference
+
+
+
#include <cuComplex.h>
+#include "cutlass/cutlass.h"
+#include <iosfwd>
+
+

Go to the source code of this file.

+ + + + +

+Classes

class  cutlass::platform::complex< T >
 
+ + + + + +

+Namespaces

 cutlass
 
 cutlass::platform
 
+ + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + +

+Functions

CUTLASS_HOST_DEVICE float const & cutlass::platform::real (cuFloatComplex const &z)
 Returns the real part of the complex number. More...
 
CUTLASS_HOST_DEVICE float & cutlass::platform::real (cuFloatComplex &z)
 Returns the real part of the complex number. More...
 
CUTLASS_HOST_DEVICE double const & cutlass::platform::real (cuDoubleComplex const &z)
 Returns the real part of the complex number. More...
 
CUTLASS_HOST_DEVICE double & cutlass::platform::real (cuDoubleComplex &z)
 Returns the real part of the complex number. More...
 
CUTLASS_HOST_DEVICE float const & cutlass::platform::imag (cuFloatComplex const &z)
 Returns the imaginary part of the complex number. More...
 
CUTLASS_HOST_DEVICE float & cutlass::platform::imag (cuFloatComplex &z)
 Returns the imaginary part of the complex number. More...
 
CUTLASS_HOST_DEVICE double const & cutlass::platform::imag (cuDoubleComplex const &z)
 Returns the imaginary part of the complex number. More...
 
CUTLASS_HOST_DEVICE double & cutlass::platform::imag (cuDoubleComplex &z)
 Returns the imaginary part of the complex number. More...
 
template<typename T >
CUTLASS_HOST_DEVICE T const & cutlass::platform::real (complex< T > const &z)
 Returns the real part of the complex number. More...
 
template<typename T >
CUTLASS_HOST_DEVICE T & cutlass::platform::real (complex< T > &z)
 Returns the real part of the complex number. More...
 
template<typename T >
CUTLASS_HOST_DEVICE T const & cutlass::platform::imag (complex< T > const &z)
 Returns the imaginary part of the complex number. More...
 
template<typename T >
CUTLASS_HOST_DEVICE T & cutlass::platform::imag (complex< T > &z)
 Returns the imaginary part of the complex number. More...
 
template<typename T >
std::ostream & cutlass::platform::operator<< (std::ostream &out, complex< T > const &z)
 
template<typename T >
CUTLASS_HOST_DEVICE bool cutlass::platform::operator== (complex< T > const &lhs, complex< T > const &rhs)
 Equality operator. More...
 
template<typename T >
CUTLASS_HOST_DEVICE bool cutlass::platform::operator!= (complex< T > const &lhs, complex< T > const &rhs)
 Inequality operator. More...
 
template<typename T >
CUTLASS_HOST_DEVICE complex< T > cutlass::platform::operator+ (complex< T > const &lhs, complex< T > const &rhs)
 Addition. More...
 
template<typename T >
CUTLASS_HOST_DEVICE complex< T > cutlass::platform::operator- (complex< T > const &lhs, complex< T > const &rhs)
 Subtraction. More...
 
template<typename T >
CUTLASS_HOST_DEVICE complex< T > cutlass::platform::operator* (complex< T > const &lhs, complex< T > const &rhs)
 Multiplication. More...
 
template<typename T >
CUTLASS_HOST_DEVICE complex< T > cutlass::platform::operator* (complex< T > const &lhs, T const &s)
 Scalar Multiplication. More...
 
template<typename T >
CUTLASS_HOST_DEVICE complex< T > cutlass::platform::operator* (T const &s, complex< T > const &rhs)
 Scalar Multiplication. More...
 
template<typename T >
CUTLASS_HOST_DEVICE complex< T > cutlass::platform::operator/ (complex< T > const &lhs, complex< T > const &rhs)
 Division. More...
 
template<typename T >
CUTLASS_HOST_DEVICE complex< T > cutlass::platform::operator/ (complex< T > const &lhs, T const &s)
 Scalar Division. More...
 
template<typename T >
CUTLASS_HOST_DEVICE complex< T > cutlass::platform::operator/ (T const &s, complex< T > const &rhs)
 Scalar divided by complex. More...
 
template<typename T >
CUTLASS_HOST_DEVICE complex< T > & cutlass::platform::operator+= (complex< T > &lhs, complex< T > const &rhs)
 Addition. More...
 
template<typename T >
CUTLASS_HOST_DEVICE complex< T > & cutlass::platform::operator-= (complex< T > &lhs, complex< T > const &rhs)
 Subtraction. More...
 
template<typename T >
CUTLASS_HOST_DEVICE complex< T > & cutlass::platform::operator*= (complex< T > &lhs, complex< T > const &rhs)
 Multiplication. More...
 
template<typename T >
CUTLASS_HOST_DEVICE complex< T > & cutlass::platform::operator*= (complex< T > &lhs, T s)
 Scalar multiplication. More...
 
template<typename T >
CUTLASS_HOST_DEVICE complex< T > & cutlass::platform::operator/= (complex< T > &lhs, complex< T > const &rhs)
 Division. More...
 
template<typename T >
CUTLASS_HOST_DEVICE T cutlass::platform::abs (complex< T > const &z)
 Returns the magnitude of the complex number. More...
 
template<typename T >
CUTLASS_HOST_DEVICE T cutlass::platform::arg (complex< T > const &z)
 Returns the magnitude of the complex number. More...
 
template<typename T >
CUTLASS_HOST_DEVICE T cutlass::platform::norm (complex< T > const &z)
 Returns the squared magnitude. More...
 
template<typename T >
CUTLASS_HOST_DEVICE complex< T > cutlass::platform::conj (complex< T > const &z)
 Returns the complex conjugate. More...
 
template<typename T >
CUTLASS_HOST_DEVICE complex< T > cutlass::platform::proj (complex< T > const &z)
 Projects the complex number z onto the Riemann sphere. More...
 
template<typename T >
CUTLASS_HOST_DEVICE complex< T > cutlass::platform::polar (T const &r, T const &theta=T())
 Returns a complex number with magnitude r and phase theta. More...
 
template<typename T >
CUTLASS_HOST_DEVICE complex< T > cutlass::platform::exp (complex< T > const &z)
 Computes the complex exponential of z. More...
 
template<typename T >
CUTLASS_HOST_DEVICE complex< T > cutlass::platform::log (complex< T > const &z)
 Computes the complex exponential of z. More...
 
template<typename T >
CUTLASS_HOST_DEVICE complex< T > cutlass::platform::log10 (complex< T > const &z)
 Computes the complex exponential of z. More...
 
template<typename T >
CUTLASS_HOST_DEVICE complex< T > cutlass::platform::sqrt (complex< T > const &z)
 Computes the square root of complex number z. More...
 
template<typename T >
CUTLASS_HOST_DEVICE complex< T > cutlass::platform::cos (complex< T > const &z)
 Computes the cosine of complex z. More...
 
template<typename T >
CUTLASS_HOST_DEVICE complex< T > cutlass::platform::sin (complex< T > const &z)
 Computes the sin of complex z. More...
 
+
+ +
+Generated on Wed Sep 19 2018 13:58:50 for Cutlass by   +doxygen + 1.8.14 +
+ + diff --git a/docs/complex_8h_source.html b/docs/complex_8h_source.html new file mode 100644 index 00000000..6270d22d --- /dev/null +++ b/docs/complex_8h_source.html @@ -0,0 +1,123 @@ + + + + + + + +Cutlass: complex.h Source File + + + + + + + + + + +
+
+ + + + + + +
+
Cutlass +
+
CUDA Templates for Linear Algebra Subroutines and Solvers
+
+
+ + + + + + + + +
+
+ + +
+ +
+ + +
+
+
+
complex.h
+
+
+Go to the documentation of this file.
1 /***************************************************************************************************
2  * Copyright (c) 2017-2018, NVIDIA CORPORATION. All rights reserved.
3  *
4  * Redistribution and use in source and binary forms, with or without modification, are permitted
5  * provided that the following conditions are met:
6  * * Redistributions of source code must retain the above copyright notice, this list of
7  * conditions and the following disclaimer.
8  * * Redistributions in binary form must reproduce the above copyright notice, this list of
9  * conditions and the following disclaimer in the documentation and/or other materials
10  * provided with the distribution.
11  * * Neither the name of the NVIDIA CORPORATION nor the names of its contributors may be used
12  * to endorse or promote products derived from this software without specific prior written
13  * permission.
14  *
15  * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND ANY EXPRESS OR
16  * IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND
17  * FITNESS FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE
18  * FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING,
19  * BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS;
20  * OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT,
21  * STRICT LIABILITY, OR TOR (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
22  * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
23  *
24  **************************************************************************************************/
25 #pragma once
26 
27 #include <cuComplex.h>
28 #include "cutlass/cutlass.h"
29 #include <iosfwd>
30 
31 namespace cutlass {
32 namespace platform {
33 
35 
36 //
37 // Accessors for CUDA complex types
38 //
39 
41 #pragma hd_warning_disable // Suppresses warnings when attempting to instantiate complex<T> with a
42  // host-only type
43 CUTLASS_HOST_DEVICE
44 float const &real(cuFloatComplex const &z) { return z.x; }
45 
47 #pragma hd_warning_disable // Suppresses warnings when attempting to instantiate complex<T> with a
48  // host-only type
49 CUTLASS_HOST_DEVICE
50 float &real(cuFloatComplex &z) { return z.x; }
51 
53 #pragma hd_warning_disable // Suppresses warnings when attempting to instantiate complex<T> with a
54  // host-only type
55 CUTLASS_HOST_DEVICE
56 double const &real(cuDoubleComplex const &z) { return z.x; }
57 
59 #pragma hd_warning_disable // Suppresses warnings when attempting to instantiate complex<T> with a
60  // host-only type
61 CUTLASS_HOST_DEVICE
62 double &real(cuDoubleComplex &z) { return z.x; }
63 
65 #pragma hd_warning_disable // Suppresses warnings when attempting to instantiate complex<T> with a
66  // host-only type
67 CUTLASS_HOST_DEVICE
68 float const &imag(cuFloatComplex const &z) { return z.y; }
69 
71 #pragma hd_warning_disable // Suppresses warnings when attempting to instantiate complex<T> with a
72  // host-only type
73 CUTLASS_HOST_DEVICE
74 float &imag(cuFloatComplex &z) { return z.y; }
75 
77 #pragma hd_warning_disable // Suppresses warnings when attempting to instantiate complex<T> with a
78  // host-only type
79 CUTLASS_HOST_DEVICE
80 double const &imag(cuDoubleComplex const &z) { return z.y; }
81 
83 #pragma hd_warning_disable // Suppresses warnings when attempting to instantiate complex<T> with a
84  // host-only type
85 CUTLASS_HOST_DEVICE
86 double &imag(cuDoubleComplex &z) { return z.y; }
87 
89 
92 template <typename T>
93 class complex {
94  public:
96  typedef T value_type;
97 
98  private:
99  //
100  // Data members
101  //
102 
104  T _real;
105 
107  T _imag;
108 
109  public:
110 //
111 // Methods
112 //
113 
115 #pragma hd_warning_disable // Suppresses warnings when attempting to instantiate complex<T> with a
116  // host-only type
117  CUTLASS_HOST_DEVICE
118  complex(T r = T(0), T i = T(0)) : _real(r), _imag(i) {}
119 
121 #pragma hd_warning_disable // Suppresses warnings when attempting to instantiate complex<T> with a
122  // host-only type
123  CUTLASS_HOST_DEVICE
124  complex(cuFloatComplex const &z) : _real(platform::real(z)), _imag(platform::imag(z)) {}
125 
127 #pragma hd_warning_disable // Suppresses warnings when attempting to instantiate complex<T> with a
128  // host-only type
129  CUTLASS_HOST_DEVICE
130  complex(cuDoubleComplex const &z) : _real(platform::real(z)), _imag(platform::imag(z)) {}
131 
133 #pragma hd_warning_disable // Suppresses warnings when attempting to instantiate complex<T> with a
134  // host-only type
135  CUTLASS_HOST_DEVICE
136  T const &real() const { return _real; }
137 
139 #pragma hd_warning_disable // Suppresses warnings when attempting to instantiate complex<T> with a
140  // host-only type
141  CUTLASS_HOST_DEVICE
142  T &real() { return _real; }
143 
145 #pragma hd_warning_disable // Suppresses warnings when attempting to instantiate complex<T> with a
146  // host-only type
147  CUTLASS_HOST_DEVICE
148  T const &imag() const { return _imag; }
149 
151 #pragma hd_warning_disable // Suppresses warnings when attempting to instantiate complex<T> with a
152  // host-only type
153  CUTLASS_HOST_DEVICE
154  T &imag() { return _imag; }
155 
157 #pragma hd_warning_disable // Suppresses warnings when attempting to instantiate complex<T> with a
158  // host-only type
159  CUTLASS_HOST_DEVICE
160  operator cuFloatComplex() const { return make_cuFloatComplex(real(), imag()); }
161 
163 #pragma hd_warning_disable // Suppresses warnings when attempting to instantiate complex<T> with a
164  // host-only type
165  CUTLASS_HOST_DEVICE
166  operator cuDoubleComplex() const { return make_cuDoubleComplex(real(), imag()); }
167 };
168 
169 //
170 // Accessors for complex template
171 //
172 
174 #pragma hd_warning_disable // Suppresses warnings when attempting to instantiate complex<T> with a
175  // host-only type
176 template <typename T>
177 CUTLASS_HOST_DEVICE T const &real(complex<T> const &z) {
178  return z.real();
179 }
180 
182 #pragma hd_warning_disable // Suppresses warnings when attempting to instantiate complex<T> with a
183  // host-only type
184 template <typename T>
185 CUTLASS_HOST_DEVICE T &real(complex<T> &z) {
186  return z.real();
187 }
188 
190 #pragma hd_warning_disable // Suppresses warnings when attempting to instantiate complex<T> with a
191  // host-only type
192 template <typename T>
193 CUTLASS_HOST_DEVICE T const &imag(complex<T> const &z) {
194  return z.imag();
195 }
196 
198 #pragma hd_warning_disable // Suppresses warnings when attempting to instantiate complex<T> with a
199  // host-only type
200 template <typename T>
201 CUTLASS_HOST_DEVICE T &imag(complex<T> &z) {
202  return z.imag();
203 }
204 
205 //
206 // Output operators
207 //
208 
209 template <typename T>
210 std::ostream &operator<<(std::ostream &out, complex<T> const &z) {
211  T _r = real(z);
212  T _i = imag(z);
213  return out << _r << "+i" << _i;
214 }
215 
216 //
217 // Non-member operators defined for complex types
218 //
219 
221 #pragma hd_warning_disable // Suppresses warnings when attempting to instantiate complex<T> with a
222  // host-only type
223 template <typename T>
224 CUTLASS_HOST_DEVICE bool operator==(complex<T> const &lhs, complex<T> const &rhs) {
225  return real(lhs) == (rhs) && imag(lhs) == imag(rhs);
226 }
227 
229 #pragma hd_warning_disable // Suppresses warnings when attempting to instantiate complex<T> with a
230  // host-only type
231 template <typename T>
232 CUTLASS_HOST_DEVICE bool operator!=(complex<T> const &lhs, complex<T> const &rhs) {
233  return !(lhs == rhs);
234 }
235 
237 #pragma hd_warning_disable // Suppresses warnings when attempting to instantiate complex<T> with a
238  // host-only type
239 template <typename T>
240 CUTLASS_HOST_DEVICE complex<T> operator+(complex<T> const &lhs, complex<T> const &rhs) {
241  return complex<T>(real(lhs) + real(rhs), imag(lhs) + imag(rhs));
242 }
243 
245 #pragma hd_warning_disable // Suppresses warnings when attempting to instantiate complex<T> with a
246  // host-only type
247 template <typename T>
248 CUTLASS_HOST_DEVICE complex<T> operator-(complex<T> const &lhs, complex<T> const &rhs) {
249  return complex<T>(real(lhs) - real(rhs), imag(lhs) - imag(rhs));
250 }
251 
253 #pragma hd_warning_disable // Suppresses warnings when attempting to instantiate complex<T> with a
254  // host-only type
255 template <typename T>
256 CUTLASS_HOST_DEVICE complex<T> operator*(complex<T> const &lhs, complex<T> const &rhs) {
257  return complex<T>(real(lhs) * real(rhs) - imag(lhs) * imag(rhs),
258  real(lhs) * imag(rhs) + imag(lhs) * real(rhs));
259 }
260 
262 #pragma hd_warning_disable // Suppresses warnings when attempting to instantiate complex<T> with a
263  // host-only type
264 template <typename T>
265 CUTLASS_HOST_DEVICE complex<T> operator*(complex<T> const &lhs, T const &s) {
266  return complex<T>(real(lhs) * s, imag(lhs) * s);
267 }
268 
270 #pragma hd_warning_disable // Suppresses warnings when attempting to instantiate complex<T> with a
271  // host-only type
272 template <typename T>
273 CUTLASS_HOST_DEVICE complex<T> operator*(T const &s, complex<T> const &rhs) {
274  return complex<T>(s * real(rhs), s * imag(rhs));
275 }
276 
278 #pragma hd_warning_disable // Suppresses warnings when attempting to instantiate complex<T> with a
279  // host-only type
280 template <typename T>
281 CUTLASS_HOST_DEVICE complex<T> operator/(complex<T> const &lhs, complex<T> const &rhs) {
282  T d = (real(rhs) * (rhs) + imag(rhs) * imag(rhs));
283 
284  return complex<T>((real(lhs) * (rhs) + imag(lhs) * imag(rhs)) / d,
285  (imag(lhs) * (rhs)-real(lhs) * imag(rhs)) / d);
286 }
287 
289 #pragma hd_warning_disable // Suppresses warnings when attempting to instantiate complex<T> with a
290  // host-only type
291 template <typename T>
292 CUTLASS_HOST_DEVICE complex<T> operator/(complex<T> const &lhs, T const &s) {
293  return complex<T>(real(lhs) / s, imag(lhs) / s);
294 }
295 
297 #pragma hd_warning_disable // Suppresses warnings when attempting to instantiate complex<T> with a
298  // host-only type
299 template <typename T>
300 CUTLASS_HOST_DEVICE complex<T> operator/(T const &s, complex<T> const &rhs) {
301  T d = (real(rhs) * (rhs) + imag(rhs) * imag(rhs));
302 
303  return complex<T>((s * (rhs)) / d, -(s * imag(rhs)) / d);
304 }
305 
307 #pragma hd_warning_disable // Suppresses warnings when attempting to instantiate complex<T> with a
308  // host-only type
309 template <typename T>
310 CUTLASS_HOST_DEVICE complex<T> &operator+=(complex<T> &lhs, complex<T> const &rhs) {
311  lhs = (lhs + rhs);
312  return lhs;
313 }
314 
316 #pragma hd_warning_disable // Suppresses warnings when attempting to instantiate complex<T> with a
317  // host-only type
318 template <typename T>
319 CUTLASS_HOST_DEVICE complex<T> &operator-=(complex<T> &lhs, complex<T> const &rhs) {
320  lhs = (lhs - rhs);
321  return lhs;
322 }
323 
325 #pragma hd_warning_disable // Suppresses warnings when attempting to instantiate complex<T> with a
326  // host-only type
327 template <typename T>
328 CUTLASS_HOST_DEVICE complex<T> &operator*=(complex<T> &lhs, complex<T> const &rhs) {
329  lhs = (lhs * rhs);
330  return lhs;
331 }
332 
334 #pragma hd_warning_disable // Suppresses warnings when attempting to instantiate complex<T> with a
335  // host-only type
336 template <typename T>
337 CUTLASS_HOST_DEVICE complex<T> &operator*=(complex<T> &lhs, T s) {
338  lhs = (lhs * s);
339  return lhs;
340 }
341 
343 #pragma hd_warning_disable // Suppresses warnings when attempting to instantiate complex<T> with a
344  // host-only type
345 template <typename T>
346 CUTLASS_HOST_DEVICE complex<T> &operator/=(complex<T> &lhs, complex<T> const &rhs) {
347  lhs = (lhs / rhs);
348  return lhs;
349 }
350 
351 //
352 // Non-member functions defined for complex numbers
353 //
354 
356 #pragma hd_warning_disable // Suppresses warnings when attempting to instantiate complex<T> with a
357  // host-only type
358 template <typename T>
359 CUTLASS_HOST_DEVICE T abs(complex<T> const &z) {
360  return sqrt(norm(z));
361 }
362 
364 #pragma hd_warning_disable // Suppresses warnings when attempting to instantiate complex<T> with a
365  // host-only type
366 template <typename T>
367 CUTLASS_HOST_DEVICE T arg(complex<T> const &z) {
368  return atan2(imag(z), real(z));
369 }
370 
372 #pragma hd_warning_disable // Suppresses warnings when attempting to instantiate complex<T> with a
373  // host-only type
374 template <typename T>
375 CUTLASS_HOST_DEVICE T norm(complex<T> const &z) {
376  return real(z) * real(z) + imag(z) * imag(z);
377 }
378 
380 #pragma hd_warning_disable // Suppresses warnings when attempting to instantiate complex<T> with a
381  // host-only type
382 template <typename T>
383 CUTLASS_HOST_DEVICE complex<T> conj(complex<T> const &z) {
384  return complex<T>(real(z), -imag(z));
385 }
386 
388 #pragma hd_warning_disable // Suppresses warnings when attempting to instantiate complex<T> with a
389  // host-only type
390 template <typename T>
391 CUTLASS_HOST_DEVICE complex<T> proj(complex<T> const &z) {
392  T d = real(z) * real(z) + imag(z) * imag(z) + T(1);
393  return complex<T>((T(2) * real(z)) / d, (T(2) * imag(z)) / d);
394 }
395 
397 #pragma hd_warning_disable // Suppresses warnings when attempting to instantiate complex<T> with a
398  // host-only type
399 template <typename T>
400 CUTLASS_HOST_DEVICE complex<T> polar(T const &r, T const &theta = T()) {
401  return complex<T>(r * cos(theta), r * sin(theta));
402 }
403 
405 #pragma hd_warning_disable // Suppresses warnings when attempting to instantiate complex<T> with a
406  // host-only type
407 template <typename T>
408 CUTLASS_HOST_DEVICE complex<T> exp(complex<T> const &z) {
409  return complex<T>(real(z) * cos(imag(z)), real(z) * sin(imag(z)));
410 }
411 
413 #pragma hd_warning_disable // Suppresses warnings when attempting to instantiate complex<T> with a
414  // host-only type
415 template <typename T>
416 CUTLASS_HOST_DEVICE complex<T> log(complex<T> const &z) {
417  return complex<T>(log(abs(z)), arg(z));
418 }
419 
421 #pragma hd_warning_disable // Suppresses warnings when attempting to instantiate complex<T> with a
422  // host-only type
423 template <typename T>
424 CUTLASS_HOST_DEVICE complex<T> log10(complex<T> const &z) {
425  return log(z) / T(log(T(10)));
426 }
427 
429 #pragma hd_warning_disable // Suppresses warnings when attempting to instantiate complex<T> with a
430  // host-only type
431 template <typename T>
432 CUTLASS_HOST_DEVICE complex<T> sqrt(complex<T> const &z) {
433  return sqrt(T(2)) / T(2) *
434  complex<T>(sqrt(sqrt(norm(z)) + real(z)),
435  (imag(z) < 0 ? T(-1) : T(1)) * sqrt(sqrt(norm(z)) - real(z)));
436 }
437 
439 #pragma hd_warning_disable // Suppresses warnings when attempting to instantiate complex<T> with a
440  // host-only type
441 template <typename T>
442 CUTLASS_HOST_DEVICE complex<T> cos(complex<T> const &z) {
443  return (exp(z) + exp(-z)) / T(2);
444 }
445 
447 #pragma hd_warning_disable // Suppresses warnings when attempting to instantiate complex<T> with a
448  // host-only type
449 template <typename T>
450 CUTLASS_HOST_DEVICE complex<T> sin(complex<T> const &z) {
451  return (exp(-z) - exp(z)) * complex<T>(T(0), T(1) / T(2));
452 }
453 
455 
456 } // namespace platform
457 } // namespace cutlass
cutlass::platform::proj
CUTLASS_HOST_DEVICE complex< T > proj(complex< T > const &z)
Projects the complex number z onto the Riemann sphere.
Definition: complex.h:391
+
cutlass
Definition: convert.h:33
+
cutlass::platform::complex::imag
CUTLASS_HOST_DEVICE T & imag()
Accesses the imaginary part of the complex number.
Definition: complex.h:154
+
cutlass::platform::operator==
CUTLASS_HOST_DEVICE bool operator==(complex< T > const &lhs, complex< T > const &rhs)
Equality operator.
Definition: complex.h:224
+
cutlass::platform::complex::imag
CUTLASS_HOST_DEVICE T const & imag() const
Accesses the imaginary part of the complex number.
Definition: complex.h:148
+
cutlass::platform::operator*
CUTLASS_HOST_DEVICE complex< T > operator*(complex< T > const &lhs, complex< T > const &rhs)
Multiplication.
Definition: complex.h:256
+
cutlass::platform::operator-=
CUTLASS_HOST_DEVICE complex< T > & operator-=(complex< T > &lhs, complex< T > const &rhs)
Subtraction.
Definition: complex.h:319
+
cutlass::platform::operator-
CUTLASS_HOST_DEVICE complex< T > operator-(complex< T > const &lhs, complex< T > const &rhs)
Subtraction.
Definition: complex.h:248
+
cutlass::platform::complex::real
CUTLASS_HOST_DEVICE T & real()
Accesses the real part of the complex number.
Definition: complex.h:142
+
cutlass::platform::real
CUTLASS_HOST_DEVICE float const & real(cuFloatComplex const &z)
Returns the real part of the complex number.
Definition: complex.h:44
+
cutlass::platform::sin
CUTLASS_HOST_DEVICE complex< T > sin(complex< T > const &z)
Computes the sin of complex z.
Definition: complex.h:450
+
cutlass::platform::complex::complex
CUTLASS_HOST_DEVICE complex(cuFloatComplex const &z)
Conversion from cuFloatComplex.
Definition: complex.h:124
+
cutlass::platform::cos
CUTLASS_HOST_DEVICE complex< T > cos(complex< T > const &z)
Computes the cosine of complex z.
Definition: complex.h:442
+
cutlass::platform::operator+
CUTLASS_HOST_DEVICE complex< T > operator+(complex< T > const &lhs, complex< T > const &rhs)
Addition.
Definition: complex.h:240
+
cutlass::platform::polar
CUTLASS_HOST_DEVICE complex< T > polar(T const &r, T const &theta=T())
Returns a complex number with magnitude r and phase theta.
Definition: complex.h:400
+
cutlass::platform::complex::real
CUTLASS_HOST_DEVICE T const & real() const
Accesses the real part of the complex number.
Definition: complex.h:136
+
cutlass::platform::operator/=
CUTLASS_HOST_DEVICE complex< T > & operator/=(complex< T > &lhs, complex< T > const &rhs)
Division.
Definition: complex.h:346
+
cutlass::platform::sqrt
CUTLASS_HOST_DEVICE complex< T > sqrt(complex< T > const &z)
Computes the square root of complex number z.
Definition: complex.h:432
+
cutlass::platform::operator+=
CUTLASS_HOST_DEVICE complex< T > & operator+=(complex< T > &lhs, complex< T > const &rhs)
Addition.
Definition: complex.h:310
+
CUTLASS_HOST_DEVICE
#define CUTLASS_HOST_DEVICE
Definition: cutlass.h:46
+
cutlass::platform::imag
CUTLASS_HOST_DEVICE float const & imag(cuFloatComplex const &z)
Returns the imaginary part of the complex number.
Definition: complex.h:68
+
cutlass::platform::exp
CUTLASS_HOST_DEVICE complex< T > exp(complex< T > const &z)
Computes the complex exponential of z.
Definition: complex.h:408
+
cutlass::platform::log10
CUTLASS_HOST_DEVICE complex< T > log10(complex< T > const &z)
Computes the complex exponential of z.
Definition: complex.h:424
+
cutlass::platform::norm
CUTLASS_HOST_DEVICE T norm(complex< T > const &z)
Returns the squared magnitude.
Definition: complex.h:375
+
cutlass::platform::operator!=
CUTLASS_HOST_DEVICE bool operator!=(complex< T > const &lhs, complex< T > const &rhs)
Inequality operator.
Definition: complex.h:232
+
cutlass::platform::abs
CUTLASS_HOST_DEVICE T abs(complex< T > const &z)
Returns the magnitude of the complex number.
Definition: complex.h:359
+
cutlass::platform::operator*=
CUTLASS_HOST_DEVICE complex< T > & operator*=(complex< T > &lhs, complex< T > const &rhs)
Multiplication.
Definition: complex.h:328
+
cutlass::platform::complex::complex
CUTLASS_HOST_DEVICE complex(cuDoubleComplex const &z)
Conversion from cuDoubleComplex.
Definition: complex.h:130
+
cutlass::platform::arg
CUTLASS_HOST_DEVICE T arg(complex< T > const &z)
Returns the magnitude of the complex number.
Definition: complex.h:367
+
cutlass::platform::complex::complex
CUTLASS_HOST_DEVICE complex(T r=T(0), T i=T(0))
Constructor.
Definition: complex.h:118
+
cutlass::platform::complex
Definition: complex.h:93
+
cutlass::platform::log
CUTLASS_HOST_DEVICE complex< T > log(complex< T > const &z)
Computes the complex exponential of z.
Definition: complex.h:416
+
cutlass::platform::complex::value_type
T value_type
Type alias for scalar type.
Definition: complex.h:96
+
cutlass.h
Basic include for CUTLASS macros.
+
cutlass::platform::operator/
CUTLASS_HOST_DEVICE complex< T > operator/(complex< T > const &lhs, complex< T > const &rhs)
Division.
Definition: complex.h:281
+
cutlass::platform::conj
CUTLASS_HOST_DEVICE complex< T > conj(complex< T > const &z)
Returns the complex conjugate.
Definition: complex.h:383
+
+ +
+Generated on Wed Sep 19 2018 13:58:50 for Cutlass by   +doxygen + 1.8.14 +
+ + diff --git a/docs/convert_8h.html b/docs/convert_8h.html index 422c5201..cd3bf4bb 100644 --- a/docs/convert_8h.html +++ b/docs/convert_8h.html @@ -82,7 +82,7 @@ $(function() {

Defines conversion operations among Fragments of different base type. More...

-
#include <cutlass/fragment.h>
+
#include "cutlass/fragment.h"

Go to the source code of this file.

@@ -103,7 +103,7 @@ Namespaces
-Generated on Thu May 3 2018 16:36:12 for Cutlass by   +Generated on Wed Sep 19 2018 13:58:50 for Cutlass by   doxygen 1.8.14
diff --git a/docs/convert_8h_source.html b/docs/convert_8h_source.html index 6e877d29..22ec9d4b 100644 --- a/docs/convert_8h_source.html +++ b/docs/convert_8h_source.html @@ -76,7 +76,7 @@ $(function() {
convert.h
-Go to the documentation of this file.
1 /***************************************************************************************************
2  * Copyright (c) 2017-2018, NVIDIA CORPORATION. All rights reserved.
3  *
4  * Redistribution and use in source and binary forms, with or without modification, are permitted
5  * provided that the following conditions are met:
6  * * Redistributions of source code must retain the above copyright notice, this list of
7  * conditions and the following disclaimer.
8  * * Redistributions in binary form must reproduce the above copyright notice, this list of
9  * conditions and the following disclaimer in the documentation and/or other materials
10  * provided with the distribution.
11  * * Neither the name of the NVIDIA CORPORATION nor the names of its contributors may be used
12  * to endorse or promote products derived from this software without specific prior written
13  * permission.
14  *
15  * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND ANY EXPRESS OR
16  * IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND
17  * FITNESS FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE
18  * FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING,
19  * BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS;
20  * OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT,
21  * STRICT LIABILITY, OR TOR (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
22  * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
23  *
24  **************************************************************************************************/
29 #pragma once
30 
31 #include <cutlass/fragment.h>
32 
33 namespace cutlass {
34 
36 
37 template <typename InputFragment_, typename OutputFragment_>
38 struct Convert {};
39 
41 
42 template <typename InputScalar_, typename OutputScalar_, int kScalars_>
43 struct Convert<Fragment<InputScalar_, kScalars_>, Fragment<OutputScalar_, kScalars_> > {
45  typedef Fragment<InputScalar_, kScalars_> InputFragment;
47  typedef Fragment<OutputScalar_, kScalars_> OutputFragment;
48 
50  CUTLASS_DEVICE Convert() {}
51 
53  CUTLASS_DEVICE void transform(InputFragment const& src, OutputFragment& dst) {
54  transform(src, 0, dst);
55  }
56 
58  template <typename Fragment_>
59  CUTLASS_DEVICE void transform(Fragment_ const& src, int offset, OutputFragment& dst) {
60  for (int i = 0; i < kScalars_; ++i) {
61  dst[i] = static_cast<OutputScalar_>(src[i + offset]);
62  }
63  }
64 };
65 
67 
68 template <typename Fragment_>
69 struct Copy {
71  typedef Fragment_ InputFragment;
73  typedef Fragment_ OutputFragment;
74 
76  CUTLASS_DEVICE Copy() {}
77 
79  CUTLASS_DEVICE void transform(Fragment_ const& src, Fragment_& dst) { transform(src, 0, dst); }
80 
82  template <typename InputFragment_>
83  CUTLASS_DEVICE void transform(InputFragment_ const& src, int offset, Fragment_& dst) {
84  if (sizeof(typename Fragment_::Element) == 8) {
85  uint64_t const* src_ptr = reinterpret_cast<uint64_t const*>(&src[offset]);
86  uint64_t* dst_ptr = reinterpret_cast<uint64_t*>(&dst[0]);
87  for (int i = 0; i < sizeof(Fragment_) / 8; ++i) {
88  dst_ptr[i] = src_ptr[i];
89  }
90  } else {
91  uint32_t const* src_ptr = reinterpret_cast<uint32_t const*>(&src[offset]);
92  uint32_t* dst_ptr = reinterpret_cast<uint32_t*>(&dst[0]);
93  for (int i = 0; i < sizeof(Fragment_) / 4; ++i) {
94  dst_ptr[i] = src_ptr[i];
95  }
96  }
97  }
98 };
99 
101 
102 } // namespace cutlass
cutlass
Definition: convert.h:33
+Go to the documentation of this file.
1 /***************************************************************************************************
2  * Copyright (c) 2017-2018, NVIDIA CORPORATION. All rights reserved.
3  *
4  * Redistribution and use in source and binary forms, with or without modification, are permitted
5  * provided that the following conditions are met:
6  * * Redistributions of source code must retain the above copyright notice, this list of
7  * conditions and the following disclaimer.
8  * * Redistributions in binary form must reproduce the above copyright notice, this list of
9  * conditions and the following disclaimer in the documentation and/or other materials
10  * provided with the distribution.
11  * * Neither the name of the NVIDIA CORPORATION nor the names of its contributors may be used
12  * to endorse or promote products derived from this software without specific prior written
13  * permission.
14  *
15  * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND ANY EXPRESS OR
16  * IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND
17  * FITNESS FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE
18  * FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING,
19  * BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS;
20  * OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT,
21  * STRICT LIABILITY, OR TOR (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
22  * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
23  *
24  **************************************************************************************************/
29 #pragma once
30 
31 #include "cutlass/fragment.h"
32 
33 namespace cutlass {
34 
36 
37 template <typename InputFragment_, typename OutputFragment_>
38 struct Convert {};
39 
41 
42 template <typename InputScalar_, typename OutputScalar_, int kScalars_>
43 struct Convert<Fragment<InputScalar_, kScalars_>, Fragment<OutputScalar_, kScalars_> > {
45  typedef Fragment<InputScalar_, kScalars_> InputFragment;
47  typedef Fragment<OutputScalar_, kScalars_> OutputFragment;
48 
50  CUTLASS_DEVICE Convert() {}
51 
53  CUTLASS_DEVICE void transform(InputFragment const& src, OutputFragment& dst) {
54  transform(src, 0, dst);
55  }
56 
58  template <typename Fragment_>
59  CUTLASS_DEVICE void transform(Fragment_ const& src, int offset, OutputFragment& dst) {
60  for (int i = 0; i < kScalars_; ++i) {
61  dst[i] = static_cast<OutputScalar_>(src[i + offset]);
62  }
63  }
64 };
65 
67 
68 template <typename Fragment_>
69 struct Copy {
71  typedef Fragment_ InputFragment;
73  typedef Fragment_ OutputFragment;
74 
76  CUTLASS_DEVICE Copy() {}
77 
79  CUTLASS_DEVICE void transform(Fragment_ const& src, Fragment_& dst) { transform(src, 0, dst); }
80 
82  template <typename InputFragment_>
83  CUTLASS_DEVICE void transform(InputFragment_ const& src, int offset, Fragment_& dst) {
84  if (sizeof(typename Fragment_::Element) == 8) {
85  uint64_t const* src_ptr = reinterpret_cast<uint64_t const*>(&src[offset]);
86  uint64_t* dst_ptr = reinterpret_cast<uint64_t*>(&dst[0]);
87  for (int i = 0; i < sizeof(Fragment_) / 8; ++i) {
88  dst_ptr[i] = src_ptr[i];
89  }
90  } else {
91  uint32_t const* src_ptr = reinterpret_cast<uint32_t const*>(&src[offset]);
92  uint32_t* dst_ptr = reinterpret_cast<uint32_t*>(&dst[0]);
93  for (int i = 0; i < sizeof(Fragment_) / 4; ++i) {
94  dst_ptr[i] = src_ptr[i];
95  }
96  }
97  }
98 };
99 
101 
102 } // namespace cutlass
cutlass
Definition: convert.h:33
cutlass::Convert< Fragment< InputScalar_, kScalars_ >, Fragment< OutputScalar_, kScalars_ > >::OutputFragment
Fragment< OutputScalar_, kScalars_ > OutputFragment
The output fragment.
Definition: convert.h:47
cutlass::Copy
Definition: convert.h:69
cutlass::Copy::transform
CUTLASS_DEVICE void transform(Fragment_ const &src, Fragment_ &dst)
Transform a fragment.
Definition: convert.h:79
@@ -94,7 +94,7 @@ $(function() {
-Generated on Thu May 3 2018 16:36:12 for Cutlass by   +Generated on Wed Sep 19 2018 13:58:50 for Cutlass by   doxygen 1.8.14
diff --git a/docs/coord_8h.html b/docs/coord_8h.html index 51650386..8bb9bea4 100644 --- a/docs/coord_8h.html +++ b/docs/coord_8h.html @@ -83,7 +83,8 @@ $(function() {

A Coord is a coordinate of arbitrary rank into a tensor or matrix. More...

-
#include <cutlass/cutlass.h>
+
#include "cutlass/cutlass.h"
+#include "cutlass/util/platform.h"

Go to the source code of this file.

@@ -92,7 +93,7 @@ Classes - +
struct  cutlass::Identity
 Describes identity elements. More...
 
struct  cutlass::Coord< N_ >
struct  cutlass::Coord< Rank_, Index_ >
 Statically-sized array specifying Coords within a tensor. More...
 
@@ -115,23 +116,14 @@ Functions - - - - - - - - - - - - + + +
CUTLASS_HOST_DEVICE Coord< 4 > cutlass::make_Coord (int _0, int _1, int _2, int _3)
 Helper to make a 4-element coordinate. More...
 
CUTLASS_HOST_DEVICE Coord< 2 > cutlass::get_Coord_hw (Coord< 3 > const &coord)
 Getter. More...
 
CUTLASS_HOST_DEVICE Coord< 2 > cutlass::get_Coord_hw (Coord< 4 > const &coord)
 Getter. More...
 
CUTLASS_HOST_DEVICE Coord< 3 > cutlass::get_Coord_hwc (Coord< 4 > const &coord)
 Getter. More...
 
CUTLASS_HOST_DEVICE Coord< 3 > cutlass::get_Coord_dhw (Coord< 4 > const &coord)
 Getter. More...
 
template<typename Shape_ >
CUTLASS_HOST_DEVICE Coord< 3 > cutlass::make_Coord_from_shape ()
 
-Generated on Thu May 3 2018 16:36:12 for Cutlass by   +Generated on Wed Sep 19 2018 13:58:50 for Cutlass by   doxygen 1.8.14
diff --git a/docs/coord_8h_source.html b/docs/coord_8h_source.html index 71ec92e1..b0e2162c 100644 --- a/docs/coord_8h_source.html +++ b/docs/coord_8h_source.html @@ -76,50 +76,54 @@ $(function() {
coord.h
-Go to the documentation of this file.
1 /***************************************************************************************************
2  * Copyright (c) 2017-2018, NVIDIA CORPORATION. All rights reserved.
3  *
4  * Redistribution and use in source and binary forms, with or without modification, are permitted
5  * provided that the following conditions are met:
6  * * Redistributions of source code must retain the above copyright notice, this list of
7  * conditions and the following disclaimer.
8  * * Redistributions in binary form must reproduce the above copyright notice, this list of
9  * conditions and the following disclaimer in the documentation and/or other materials
10  * provided with the distribution.
11  * * Neither the name of the NVIDIA CORPORATION nor the names of its contributors may be used
12  * to endorse or promote products derived from this software without specific prior written
13  * permission.
14  *
15  * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND ANY EXPRESS OR
16  * IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND
17  * FITNESS FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE
18  * FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING,
19  * BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS;
20  * OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT,
21  * STRICT LIABILITY, OR TOR (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
22  * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
23  *
24  **************************************************************************************************/
29 #pragma once
30 
31 #include <cutlass/cutlass.h>
32 
33 namespace cutlass {
34 
36 
38 struct Identity {
41  enum Kind { Additive = 0, Multiplicative = 1 };
42 };
43 
45 
47 template <int N_>
48 struct Coord {
49  //
50  // Type and constant definitions
51  //
52 
53  static int const N = N_;
54 
55  //
56  // Data members
57  //
58 
60  int idx[N];
61 
62  //
63  // Methods
64  //
65 
67  CUTLASS_HOST_DEVICE
68  Coord(int value = 0) {
69  for (int i = 0; i < N; ++i) {
70  idx[i] = value;
71  }
72  }
73 
75  CUTLASS_HOST_DEVICE
76  Coord(int _idx[]) {
77  for (int i = 0; i < N; ++i) {
78  idx[i] = _idx[i];
79  }
80  }
81 
83  CUTLASS_HOST_DEVICE
84  Coord operator+(Coord const& b) const {
85  Coord c;
86  for (int i = 0; i < N; ++i) {
87  c.idx[i] = idx[i] + b.idx[i];
88  }
89  return c;
90  }
91 
93  CUTLASS_HOST_DEVICE
94  Coord operator-(Coord const& b) const {
95  Coord c;
96  for (int i = 0; i < N; ++i) {
97  c.idx[i] = idx[i] - b.idx[i];
98  }
99  return c;
100  }
101 
103  CUTLASS_HOST_DEVICE
104  Coord operator*(Coord const& b) const {
105  Coord c;
106  for (int i = 0; i < N; ++i) {
107  c.idx[i] = idx[i] * b.idx[i];
108  }
109  return c;
110  }
111 
113  CUTLASS_HOST_DEVICE
114  Coord operator/(Coord const& b) const {
115  Coord c;
116  for (int i = 0; i < N; ++i) {
117  c.idx[i] = idx[i] / b.idx[i];
118  }
119  return c;
120  }
121 
123  CUTLASS_HOST_DEVICE
124  Coord& operator+=(Coord const& b) {
125  for (int i = 0; i < N; ++i) {
126  idx[i] += b.idx[i];
127  }
128  return *this;
129  }
130 
132  CUTLASS_HOST_DEVICE
133  Coord& operator-=(Coord const& b) {
134  for (int i = 0; i < N; ++i) {
135  idx[i] -= b.idx[i];
136  }
137  return *this;
138  }
139 
141  CUTLASS_HOST_DEVICE
142  Coord& operator*=(Coord const& b) {
143  for (int i = 0; i < N; ++i) {
144  idx[i] *= b.idx[i];
145  }
146  return *this;
147  }
148 
150  CUTLASS_HOST_DEVICE
151  Coord& operator/=(Coord const& b) {
152  for (int i = 0; i < N; ++i) {
153  idx[i] /= b.idx[i];
154  }
155  return *this;
156  }
157 
159  CUTLASS_HOST_DEVICE int& operator[](int dim) { return idx[dim]; }
160 
162  CUTLASS_HOST_DEVICE int const& operator[](int dim) const { return idx[dim]; }
163 
165  template <typename T>
166  CUTLASS_HOST_DEVICE T dot(Coord const& b, T sum) const {
167  for (int i = 0; i < N; ++i) {
168  sum += idx[i] * b.idx[i];
169  }
170  return sum;
171  }
172 
174  template <typename T>
175  CUTLASS_HOST_DEVICE T dot(Coord const& b) const {
176  T sum = T(0);
177  for (int i = 0; i < N; ++i) {
178  sum += idx[i] * b.idx[i];
179  }
180  return sum;
181  }
182 
184  template <int Dim>
185  CUTLASS_HOST_DEVICE int& at() {
186  return idx[Dim];
187  }
188 
190  CUTLASS_HOST_DEVICE
191  int& at(int dim) { return idx[dim]; }
192 
194  template <int Dim>
195  CUTLASS_HOST_DEVICE int const& at() const {
196  return idx[Dim];
197  }
198 
200  CUTLASS_HOST_DEVICE
201  int const& at(int dim) const { return idx[dim]; }
202 
204  CUTLASS_HOST_DEVICE
205  bool operator==(Coord<N> const& b) const {
206  bool equal = true;
207  for (int i = 0; equal && i < N; ++i) {
208  equal = (idx[i] == b.idx[i]);
209  }
210  return equal;
211  }
212 
214  CUTLASS_HOST_DEVICE
215  bool operator!=(Coord<N> const& b) const { return !(*this == b); }
216 
218  CUTLASS_HOST_DEVICE
219  Coord& clamp(Coord<N> const& max, Coord<N> const& min = Coord<N>()) {
220  for (int i = 0; i < N; ++i) {
221  idx[i] = __NV_STD_MAX(__NV_STD_MIN(idx[i], max.idx[i]), min.idx[i]);
222  }
223  return *this;
224  }
225 
227  CUTLASS_HOST_DEVICE
228  int count() const {
229  int product = idx[0];
230  for (int i = 1; i < N; ++i) {
231  product *= idx[i];
232  }
233  return product;
234  }
235 };
236 
238 
240 CUTLASS_HOST_DEVICE
241 Coord<1> make_Coord(int _0) {
242  int values[1] = {_0};
243  return Coord<1>(values);
244 }
245 
247 CUTLASS_HOST_DEVICE
248 Coord<2> make_Coord(int _0, int _1) {
249  int values[2] = {_0, _1};
250  return Coord<2>(values);
251 }
252 
254 CUTLASS_HOST_DEVICE
255 Coord<3> make_Coord(int _0, int _1, int _2) {
256  int values[3] = {_0, _1, _2};
257  return Coord<3>(values);
258 }
259 
261 CUTLASS_HOST_DEVICE
262 Coord<4> make_Coord(int _0, int _1, int _2, int _3) {
263  int values[4] = {_0, _1, _2, _3};
264  return Coord<4>(values);
265 }
266 
268 
270 CUTLASS_HOST_DEVICE
271 Coord<2> get_Coord_hw(Coord<3> const& coord) { return make_Coord(coord[1], coord[2]); }
272 
274 CUTLASS_HOST_DEVICE
275 Coord<2> get_Coord_hw(Coord<4> const& coord) { return make_Coord(coord[1], coord[2]); }
276 
278 CUTLASS_HOST_DEVICE
279 Coord<3> get_Coord_hwc(Coord<4> const& coord) { return make_Coord(coord[1], coord[2], coord[3]); }
280 
282 CUTLASS_HOST_DEVICE
283 Coord<3> get_Coord_dhw(Coord<4> const& coord) { return make_Coord(coord[0], coord[1], coord[2]); }
284 
286 
287 } // namespace cutlass
cutlass::Coord::operator[]
CUTLASS_HOST_DEVICE int const & operator[](int dim) const
Member access operator.
Definition: coord.h:162
-
cutlass::Coord::count
CUTLASS_HOST_DEVICE int count() const
Returns the product of all elements.
Definition: coord.h:228
-
cutlass::Identity
Describes identity elements.
Definition: coord.h:38
-
cutlass::platform::max
CUTLASS_HOST_DEVICE constexpr const T & max(const T &a, const T &b)
std::max
Definition: platform.h:207
+Go to the documentation of this file.
1 /***************************************************************************************************
2  * Copyright (c) 2017-2018, NVIDIA CORPORATION. All rights reserved.
3  *
4  * Redistribution and use in source and binary forms, with or without modification, are permitted
5  * provided that the following conditions are met:
6  * * Redistributions of source code must retain the above copyright notice, this list of
7  * conditions and the following disclaimer.
8  * * Redistributions in binary form must reproduce the above copyright notice, this list of
9  * conditions and the following disclaimer in the documentation and/or other materials
10  * provided with the distribution.
11  * * Neither the name of the NVIDIA CORPORATION nor the names of its contributors may be used
12  * to endorse or promote products derived from this software without specific prior written
13  * permission.
14  *
15  * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND ANY EXPRESS OR
16  * IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND
17  * FITNESS FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE
18  * FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING,
19  * BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS;
20  * OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT,
21  * STRICT LIABILITY, OR TOR (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
22  * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
23  *
24  **************************************************************************************************/
29 #pragma once
30 
31 #include "cutlass/cutlass.h"
32 #include "cutlass/util/platform.h"
33 
34 namespace cutlass {
35 
37 
39 struct Identity {
42  enum Kind { Additive = 0, Multiplicative = 1 };
43 };
44 
46 
48 template <int Rank_, typename Index_ = int>
49 struct Coord {
50  //
51  // Type and constant definitions
52  //
53 
55  static int const kRank = Rank_;
56 
58  static int const N = Rank_;
59 
61  typedef Index_ Index;
62 
63  //
64  // Data members
65  //
66 
68  Index idx[kRank];
69 
70  //
71  // Methods
72  //
73 
75  CUTLASS_HOST_DEVICE
76  Coord(Index value = 0) {
77  for (int i = 0; i < kRank; ++i) {
78  idx[i] = value;
79  }
80  }
81 
83  CUTLASS_HOST_DEVICE
84  Coord(Index _idx[]) {
85  for (int i = 0; i < kRank; ++i) {
86  idx[i] = _idx[i];
87  }
88  }
89 
91  CUTLASS_HOST_DEVICE
92  Coord(Coord<kRank> const &coord) {
93  for (int i = 0; i < kRank; ++i) {
94  idx[i] = coord[i];
95  }
96  }
97 
100  template <int Slice>
101  CUTLASS_HOST_DEVICE
102  Coord<Slice> slice(int start = 0, Index identity = 0) const {
103  Coord<Slice> result;
104  for (int i = 0; i < Slice; ++i) {
105  if (i + start < kRank) {
106  slice[i] = idx[i + start];
107  }
108  else {
109  slice[i] = identity;
110  }
111  }
112  return result;
113  }
114 
116  CUTLASS_HOST_DEVICE
117  operator bool() const {
118  for (int i = 0; i < kRank; ++i) {
119  if (idx[i]) {
120  return true;
121  }
122  }
123  return false;
124  }
125 
127  CUTLASS_HOST_DEVICE
128  bool operator!() const {
129  for (int i = 0; i < kRank; ++i) {
130  if (idx[i]) {
131  return false;
132  }
133  }
134  return true;
135  }
136 
138  CUTLASS_HOST_DEVICE
139  Coord operator+(Coord const& b) const {
140  Coord c;
141  for (int i = 0; i < kRank; ++i) {
142  c.idx[i] = idx[i] + b.idx[i];
143  }
144  return c;
145  }
146 
148  CUTLASS_HOST_DEVICE
149  Coord operator-(Coord const& b) const {
150  Coord c;
151  for (int i = 0; i < kRank; ++i) {
152  c.idx[i] = idx[i] - b.idx[i];
153  }
154  return c;
155  }
156 
158  CUTLASS_HOST_DEVICE
159  Coord operator*(Coord const& b) const {
160  Coord c;
161  for (int i = 0; i < kRank; ++i) {
162  c.idx[i] = idx[i] * b.idx[i];
163  }
164  return c;
165  }
166 
168  CUTLASS_HOST_DEVICE
169  Coord operator/(Coord const& b) const {
170  Coord c;
171  for (int i = 0; i < kRank; ++i) {
172  c.idx[i] = idx[i] / b.idx[i];
173  }
174  return c;
175  }
176 
178  CUTLASS_HOST_DEVICE
179  Coord& operator+=(Coord const& b) {
180  for (int i = 0; i < kRank; ++i) {
181  idx[i] += b.idx[i];
182  }
183  return *this;
184  }
185 
187  CUTLASS_HOST_DEVICE
188  Coord& operator-=(Coord const& b) {
189  for (int i = 0; i < kRank; ++i) {
190  idx[i] -= b.idx[i];
191  }
192  return *this;
193  }
194 
196  CUTLASS_HOST_DEVICE
197  Coord& operator*=(Coord const& b) {
198  for (int i = 0; i < kRank; ++i) {
199  idx[i] *= b.idx[i];
200  }
201  return *this;
202  }
203 
205  CUTLASS_HOST_DEVICE
206  Coord& operator/=(Coord const& b) {
207  for (int i = 0; i < kRank; ++i) {
208  idx[i] /= b.idx[i];
209  }
210  return *this;
211  }
212 
214  CUTLASS_HOST_DEVICE Index& operator[](int dim) { return idx[dim]; }
215 
217  CUTLASS_HOST_DEVICE Index const& operator[](int dim) const { return idx[dim]; }
218 
220  template <typename T>
221  CUTLASS_HOST_DEVICE T dot(Coord const& b, T sum) const {
222  for (int i = 0; i < kRank; ++i) {
223  sum += idx[i] * b.idx[i];
224  }
225  return sum;
226  }
227 
229  template <typename T>
230  CUTLASS_HOST_DEVICE T dot(Coord const& b) const {
231  T sum = T(0);
232  for (int i = 0; i < kRank; ++i) {
233  sum += idx[i] * b.idx[i];
234  }
235  return sum;
236  }
237 
239  template <int Dim>
240  CUTLASS_HOST_DEVICE Index& at() {
241  return idx[Dim];
242  }
243 
245  CUTLASS_HOST_DEVICE
246  Index& at(int dim) { return idx[dim]; }
247 
249  template <int Dim>
250  CUTLASS_HOST_DEVICE Index const& at() const {
251  return idx[Dim];
252  }
253 
255  CUTLASS_HOST_DEVICE
256  Index const& at(int dim) const { return idx[dim]; }
257 
259  CUTLASS_HOST_DEVICE
260  bool operator==(Coord<kRank> const& b) const {
261  bool equal = true;
262  for (int i = 0; equal && i < kRank; ++i) {
263  equal = (idx[i] == b.idx[i]);
264  }
265  return equal;
266  }
267 
269  CUTLASS_HOST_DEVICE
270  bool operator!=(Coord<kRank> const& b) const { return !(*this == b); }
271 
273  CUTLASS_HOST_DEVICE
274  Coord& clamp(Coord<kRank> const& max, Coord<kRank> const& min = Coord<kRank>()) {
275  for (int i = 0; i < kRank; ++i) {
276  idx[i] = __NV_STD_MAX(__NV_STD_MIN(idx[i], max.idx[i]), min.idx[i]);
277  }
278  return *this;
279  }
280 
282  CUTLASS_HOST_DEVICE
283  Index count() const {
284  Index product = idx[0];
285  for (int i = 1; i < kRank; ++i) {
286  product *= idx[i];
287  }
288  return product;
289  }
290 
292  CUTLASS_HOST_DEVICE
293  bool operator<(Coord<kRank> const &b) const {
294  for (int i = 0; i < kRank; ++i) {
295  if (!(idx[i] < b[i])) {
296  return false;
297  }
298  }
299  return true;
300  }
301 
303  CUTLASS_HOST_DEVICE
304  bool operator<=(Coord<kRank> const &b) const {
305  for (int i = 0; i < kRank; ++i) {
306  if (!(idx[i] <= b[i])) {
307  return false;
308  }
309  }
310  return true;
311  }
312 };
313 
315 
317 CUTLASS_HOST_DEVICE
318 Coord<1> make_Coord(int _0) {
319  int values[1] = {_0};
320  return Coord<1>(values);
321 }
322 
324 CUTLASS_HOST_DEVICE
325 Coord<2> make_Coord(int _0, int _1) {
326  int values[2] = {_0, _1};
327  return Coord<2>(values);
328 }
329 
331 CUTLASS_HOST_DEVICE
332 Coord<3> make_Coord(int _0, int _1, int _2) {
333  int values[3] = {_0, _1, _2};
334  return Coord<3>(values);
335 }
336 
338 CUTLASS_HOST_DEVICE
339 Coord<4> make_Coord(int _0, int _1, int _2, int _3) {
340  int values[4] = {_0, _1, _2, _3};
341  return Coord<4>(values);
342 }
343 
345 
346 template <typename Shape_>
347 CUTLASS_HOST_DEVICE Coord<3> make_Coord_from_shape() {
348  return make_Coord(Shape_::kD, Shape_::kH, Shape_::kW);
349 }
350 
352 
353 } // namespace cutlass
cutlass::Identity
Describes identity elements.
Definition: coord.h:39
+
cutlass::platform::max
CUTLASS_HOST_DEVICE constexpr const T & max(const T &a, const T &b)
std::max
Definition: platform.h:215
cutlass
Definition: convert.h:33
-
cutlass::Coord::operator==
CUTLASS_HOST_DEVICE bool operator==(Coord< N > const &b) const
Determines if two Coord<> objects are equal.
Definition: coord.h:205
-
cutlass::Coord::operator+=
CUTLASS_HOST_DEVICE Coord & operator+=(Coord const &b)
In-place addition.
Definition: coord.h:124
-
cutlass::Coord::operator!=
CUTLASS_HOST_DEVICE bool operator!=(Coord< N > const &b) const
Not equal.
Definition: coord.h:215
-
cutlass::make_Coord
CUTLASS_HOST_DEVICE Coord< 1 > make_Coord(int _0)
Helper to make a 2-element coordinate.
Definition: coord.h:241
-
cutlass::get_Coord_hwc
CUTLASS_HOST_DEVICE Coord< 3 > get_Coord_hwc(Coord< 4 > const &coord)
Getter.
Definition: coord.h:279
-
cutlass::get_Coord_dhw
CUTLASS_HOST_DEVICE Coord< 3 > get_Coord_dhw(Coord< 4 > const &coord)
Getter.
Definition: coord.h:283
-
cutlass::Coord::clamp
CUTLASS_HOST_DEVICE Coord & clamp(Coord< N > const &max, Coord< N > const &min=Coord< N >())
Clamps a coordinate to a range specified by maximum and minimum values.
Definition: coord.h:219
-
cutlass::Coord::at
CUTLASS_HOST_DEVICE int const & at() const
Gets the index of a given Coord element.
Definition: coord.h:195
-
cutlass::Coord::operator/
CUTLASS_HOST_DEVICE Coord operator/(Coord const &b) const
Element-wise division.
Definition: coord.h:114
-
cutlass::Identity::Kind
Kind
Definition: coord.h:41
-
cutlass::Coord::dot
CUTLASS_HOST_DEVICE T dot(Coord const &b, T sum) const
Computes the dot product of two Coord instances.
Definition: coord.h:166
-
cutlass::Coord::Coord
CUTLASS_HOST_DEVICE Coord(int _idx[])
Constructs from an array of integers.
Definition: coord.h:76
-
__NV_STD_MAX
#define __NV_STD_MAX(a, b)
Select maximum(a, b)
Definition: platform.h:155
-
cutlass::Coord::at
CUTLASS_HOST_DEVICE int & at(int dim)
Access via index; may limit unrolling potential.
Definition: coord.h:191
-
cutlass::Coord::operator[]
CUTLASS_HOST_DEVICE int & operator[](int dim)
Member access operator.
Definition: coord.h:159
-
cutlass::Coord::operator-=
CUTLASS_HOST_DEVICE Coord & operator-=(Coord const &b)
In-place subtraction.
Definition: coord.h:133
-
cutlass::Coord::operator*
CUTLASS_HOST_DEVICE Coord operator*(Coord const &b) const
Element-wise multiplication.
Definition: coord.h:104
-
cutlass::Coord::Coord
CUTLASS_HOST_DEVICE Coord(int value=0)
Default ctor initializes uniformly.
Definition: coord.h:68
-
cutlass::get_Coord_hw
CUTLASS_HOST_DEVICE Coord< 2 > get_Coord_hw(Coord< 3 > const &coord)
Getter.
Definition: coord.h:271
-
cutlass::Coord::N
static int const N
Definition: coord.h:53
-
__NV_STD_MIN
#define __NV_STD_MIN(a, b)
Select minimum(a, b)
Definition: platform.h:160
-
cutlass::Coord::dot
CUTLASS_HOST_DEVICE T dot(Coord const &b) const
Computes the dot product of two Coord instances.
Definition: coord.h:175
-
cutlass::Coord::operator-
CUTLASS_HOST_DEVICE Coord operator-(Coord const &b) const
Element-wise subtraction.
Definition: coord.h:94
+
cutlass::Coord::operator-
CUTLASS_HOST_DEVICE Coord operator-(Coord const &b) const
Element-wise subtraction.
Definition: coord.h:149
+
cutlass::Coord::at
CUTLASS_HOST_DEVICE Index const & at(int dim) const
Access via index; may limit unrolling potential.
Definition: coord.h:256
+
cutlass::Coord::operator[]
CUTLASS_HOST_DEVICE Index const & operator[](int dim) const
Member access operator.
Definition: coord.h:217
+
cutlass::Coord::operator/
CUTLASS_HOST_DEVICE Coord operator/(Coord const &b) const
Element-wise division.
Definition: coord.h:169
+
cutlass::Coord::operator[]
CUTLASS_HOST_DEVICE Index & operator[](int dim)
Member access operator.
Definition: coord.h:214
+
cutlass::make_Coord
CUTLASS_HOST_DEVICE Coord< 1 > make_Coord(int _0)
Helper to make a 2-element coordinate.
Definition: coord.h:318
+
cutlass::Coord::kRank
static int const kRank
Number of elements in Coord.
Definition: coord.h:55
+
cutlass::Coord::Index
Index_ Index
Index type used to store elements.
Definition: coord.h:61
+
cutlass::Coord::operator*=
CUTLASS_HOST_DEVICE Coord & operator*=(Coord const &b)
In-place multiplication.
Definition: coord.h:197
+
cutlass::Coord::at
CUTLASS_HOST_DEVICE Index & at(int dim)
Access via index; may limit unrolling potential.
Definition: coord.h:246
+
platform.h
C++ features that may be otherwise unimplemented for CUDA device functions.
+
cutlass::Coord::count
CUTLASS_HOST_DEVICE Index count() const
Returns the product of all elements.
Definition: coord.h:283
+
cutlass::Coord::operator*
CUTLASS_HOST_DEVICE Coord operator*(Coord const &b) const
Element-wise multiplication.
Definition: coord.h:159
+
cutlass::Identity::Kind
Kind
Definition: coord.h:42
+
cutlass::make_Coord_from_shape
CUTLASS_HOST_DEVICE Coord< 3 > make_Coord_from_shape()
Definition: coord.h:347
+
cutlass::Coord::operator==
CUTLASS_HOST_DEVICE bool operator==(Coord< kRank > const &b) const
Determines if two Coord<> objects are equal.
Definition: coord.h:260
+
cutlass::Coord::N
static int const N
Number of elements in Coord, aliased for compatibility.
Definition: coord.h:58
+
__NV_STD_MAX
#define __NV_STD_MAX(a, b)
Select maximum(a, b)
Definition: platform.h:163
+
cutlass::Coord::idx
Index idx[kRank]
Indices.
Definition: coord.h:68
+
__NV_STD_MIN
#define __NV_STD_MIN(a, b)
Select minimum(a, b)
Definition: platform.h:168
+
cutlass::Coord::operator-=
CUTLASS_HOST_DEVICE Coord & operator-=(Coord const &b)
In-place subtraction.
Definition: coord.h:188
+
cutlass::Coord::operator+=
CUTLASS_HOST_DEVICE Coord & operator+=(Coord const &b)
In-place addition.
Definition: coord.h:179
CUTLASS_HOST_DEVICE
#define CUTLASS_HOST_DEVICE
Definition: cutlass.h:46
-
cutlass::platform::min
CUTLASS_HOST_DEVICE constexpr const T & min(const T &a, const T &b)
std::min
Definition: platform.h:201
-
cutlass::Identity::Multiplicative
Definition: coord.h:41
-
cutlass::Coord
Statically-sized array specifying Coords within a tensor.
Definition: coord.h:48
-
cutlass::Coord::at
CUTLASS_HOST_DEVICE int & at()
Gets the index of a given Coord element.
Definition: coord.h:185
-
cutlass::Coord::idx
int idx[N]
Indices.
Definition: coord.h:60
-
cutlass::Identity::Additive
Definition: coord.h:41
-
cutlass::Coord::at
CUTLASS_HOST_DEVICE int const & at(int dim) const
Access via index; may limit unrolling potential.
Definition: coord.h:201
+
cutlass::Coord::operator!=
CUTLASS_HOST_DEVICE bool operator!=(Coord< kRank > const &b) const
Not equal.
Definition: coord.h:270
+
cutlass::platform::min
CUTLASS_HOST_DEVICE constexpr const T & min(const T &a, const T &b)
std::min
Definition: platform.h:209
+
cutlass::Coord::at
CUTLASS_HOST_DEVICE Index & at()
Gets the index of a given Coord element.
Definition: coord.h:240
+
cutlass::Coord::operator/=
CUTLASS_HOST_DEVICE Coord & operator/=(Coord const &b)
In-place division.
Definition: coord.h:206
+
cutlass::Identity::Multiplicative
Definition: coord.h:42
+
cutlass::Coord::slice
CUTLASS_HOST_DEVICE Coord< Slice > slice(int start=0, Index identity=0) const
Definition: coord.h:102
+
cutlass::Coord
Statically-sized array specifying Coords within a tensor.
Definition: coord.h:49
+
cutlass::Coord::at
CUTLASS_HOST_DEVICE Index const & at() const
Gets the index of a given Coord element.
Definition: coord.h:250
+
cutlass::Coord::dot
CUTLASS_HOST_DEVICE T dot(Coord const &b, T sum) const
Computes the dot product of two Coord instances.
Definition: coord.h:221
+
cutlass::Coord::Coord
CUTLASS_HOST_DEVICE Coord(Index value=0)
Default ctor initializes uniformly.
Definition: coord.h:76
+
cutlass::Identity::Additive
Definition: coord.h:42
+
cutlass::Coord::clamp
CUTLASS_HOST_DEVICE Coord & clamp(Coord< kRank > const &max, Coord< kRank > const &min=Coord< kRank >())
Clamps a coordinate to a range specified by maximum and minimum values.
Definition: coord.h:274
+
cutlass::Coord::Coord
CUTLASS_HOST_DEVICE Coord(Index _idx[])
Constructs from an array of integers.
Definition: coord.h:84
+
cutlass::Coord::dot
CUTLASS_HOST_DEVICE T dot(Coord const &b) const
Computes the dot product of two Coord instances.
Definition: coord.h:230
+
cutlass::Coord::operator+
CUTLASS_HOST_DEVICE Coord operator+(Coord const &b) const
Element-wise addition.
Definition: coord.h:139
cutlass.h
Basic include for CUTLASS macros.
-
cutlass::Coord::operator*=
CUTLASS_HOST_DEVICE Coord & operator*=(Coord const &b)
In-place multiplication.
Definition: coord.h:142
-
cutlass::Coord::operator+
CUTLASS_HOST_DEVICE Coord operator+(Coord const &b) const
Element-wise addition.
Definition: coord.h:84
-
cutlass::Coord::operator/=
CUTLASS_HOST_DEVICE Coord & operator/=(Coord const &b)
In-place division.
Definition: coord.h:151
+
cutlass::Coord::Coord
CUTLASS_HOST_DEVICE Coord(Coord< kRank > const &coord)
Constructs from an array of integers.
Definition: coord.h:92
+
cutlass::Coord::operator!
CUTLASS_HOST_DEVICE bool operator!() const
Returns true if Coord is uniformly zero.
Definition: coord.h:128
-Generated on Thu May 3 2018 16:36:12 for Cutlass by   +Generated on Wed Sep 19 2018 13:58:50 for Cutlass by   doxygen 1.8.14
diff --git a/docs/core__io_8h.html b/docs/core__io_8h.html index d71c3971..2f50d785 100644 --- a/docs/core__io_8h.html +++ b/docs/core__io_8h.html @@ -73,6 +73,8 @@ $(function() {
+Classes | +Namespaces | Functions
core_io.h File Reference
@@ -83,51 +85,56 @@ $(function() { More...

#include <iosfwd>
#include <typeinfo>
-#include <cutlass/coord.h>
+#include "cutlass/coord.h"
+#include "cutlass/vector.h"

Go to the source code of this file.

+ + + + +

+Classes

struct  cutlass::ScalarIO< T >
 Helper to enable formatted printing of CUTLASS scalar types to an ostream. More...
 
+ + + +

+Namespaces

 cutlass
 
- - - + + + + + + + + + + + + + + + + + + + + + + + + + + +

Functions

template<int Rank>
std::ostream & operator<< (std::ostream &out, cutlass::Coord< Rank > const &coord)
 
template<int Rank>
std::ostream & cutlass::operator<< (std::ostream &out, Coord< Rank > const &coord)
 
template<typename T >
std::ostream & cutlass::operator<< (std::ostream &out, ScalarIO< T > const &scalar)
 Default printing to ostream. More...
 
template<>
std::ostream & cutlass::operator<< (std::ostream &out, ScalarIO< int8_t > const &scalar)
 Printing to ostream of int8_t as integer rather than character. More...
 
template<>
std::ostream & cutlass::operator<< (std::ostream &out, ScalarIO< uint8_t > const &scalar)
 Printing to ostream of uint8_t as integer rather than character. More...
 
template<>
std::ostream & cutlass::operator<< (std::ostream &out, ScalarIO< cutlass::Vector< cutlass::bin1_t, 32 > > const &scalar)
 Printing to ostream of vector of 1b elements. More...
 
template<>
std::ostream & cutlass::operator<< (std::ostream &out, ScalarIO< cutlass::Vector< cutlass::int4_t, 8 > > const &scalar)
 Printing to ostream of vector of 4b signed integer elements. More...
 
template<>
std::ostream & cutlass::operator<< (std::ostream &out, ScalarIO< cutlass::Vector< cutlass::uint4_t, 8 > > const &scalar)
 Printing to ostream of vector of 4b unsigned integer elements. More...
 
-

Function Documentation

- -

◆ operator<<()

- -
-
-
-template<int Rank>
- - - - - - - - - - - - - - - - - - -
std::ostream& operator<< (std::ostream & out,
cutlass::Coord< Rank > const & coord 
)
-
- -
-
-Generated on Thu May 3 2018 16:36:12 for Cutlass by   +Generated on Wed Sep 19 2018 13:58:50 for Cutlass by   doxygen 1.8.14
diff --git a/docs/core__io_8h_source.html b/docs/core__io_8h_source.html index 7c076c94..21b79011 100644 --- a/docs/core__io_8h_source.html +++ b/docs/core__io_8h_source.html @@ -76,11 +76,19 @@ $(function() {
core_io.h
-Go to the documentation of this file.
1 /***************************************************************************************************
2  * Copyright (c) 2017-2018, NVIDIA CORPORATION. All rights reserved.
3  *
4  * Redistribution and use in source and binary forms, with or without modification, are permitted
5  * provided that the following conditions are met:
6  * * Redistributions of source code must retain the above copyright notice, this list of
7  * conditions and the following disclaimer.
8  * * Redistributions in binary form must reproduce the above copyright notice, this list of
9  * conditions and the following disclaimer in the documentation and/or other materials
10  * provided with the distribution.
11  * * Neither the name of the NVIDIA CORPORATION nor the names of its contributors may be used
12  * to endorse or promote products derived from this software without specific prior written
13  * permission.
14  *
15  * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND ANY EXPRESS OR
16  * IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND
17  * FITNESS FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE
18  * FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING,
19  * BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS;
20  * OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT,
21  * STRICT LIABILITY, OR TOR (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
22  * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
23  *
24  **************************************************************************************************/
25 #pragma once
26 
31 #pragma once
32 
33 #include <iosfwd>
34 #include <typeinfo>
35 
36 #include <cutlass/coord.h>
37 
38 template <int Rank>
39 std::ostream& operator<<(std::ostream& out, cutlass::Coord<Rank> const& coord) {
40  for (int i = 0; i < Rank; ++i) {
41  out << (i ? ", " : "") << coord.idx[i];
42  }
43  return out;
44 }
coord.h
A Coord is a coordinate of arbitrary rank into a tensor or matrix.
+Go to the documentation of this file.
1 /***************************************************************************************************
2  * Copyright (c) 2017-2018, NVIDIA CORPORATION. All rights reserved.
3  *
4  * Redistribution and use in source and binary forms, with or without modification, are permitted
5  * provided that the following conditions are met:
6  * * Redistributions of source code must retain the above copyright notice, this list of
7  * conditions and the following disclaimer.
8  * * Redistributions in binary form must reproduce the above copyright notice, this list of
9  * conditions and the following disclaimer in the documentation and/or other materials
10  * provided with the distribution.
11  * * Neither the name of the NVIDIA CORPORATION nor the names of its contributors may be used
12  * to endorse or promote products derived from this software without specific prior written
13  * permission.
14  *
15  * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND ANY EXPRESS OR
16  * IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND
17  * FITNESS FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE
18  * FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING,
19  * BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS;
20  * OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT,
21  * STRICT LIABILITY, OR TOR (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
22  * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
23  *
24  **************************************************************************************************/
29 #pragma once
30 
31 #include <iosfwd>
32 #include <typeinfo>
33 
34 #include "cutlass/coord.h"
35 #include "cutlass/vector.h"
36 
37 namespace cutlass {
38 
40 
41 template <int Rank>
42 std::ostream& operator<<(std::ostream& out, Coord<Rank> const& coord) {
43  for (int i = 0; i < Rank; ++i) {
44  out << (i ? ", " : "") << coord.idx[i];
45  }
46  return out;
47 }
48 
50 
52 template <typename T>
53 struct ScalarIO {
54 
56  T value;
57 
59  ScalarIO() { }
60 
62  ScalarIO(T value): value(value) {}
63 };
64 
66 
68 template <typename T>
69 inline std::ostream &operator<<(std::ostream &out, ScalarIO<T> const &scalar) {
70  return out << scalar.value;
71 }
72 
74 template <>
75 inline std::ostream &operator<<(std::ostream &out, ScalarIO<int8_t> const &scalar) {
76  return out << int(scalar.value);
77 }
78 
80 template <>
81 inline std::ostream &operator<<(std::ostream &out, ScalarIO<uint8_t> const &scalar) {
82  return out << unsigned(scalar.value);
83 }
84 
86 template <>
87 inline std::ostream &operator<<(
88  std::ostream &out,
89  ScalarIO<cutlass::Vector<cutlass::bin1_t, 32> > const &scalar) {
90 
91  for (int i = 0; i < 32; i++) {
92  out << int(scalar.value[i]);
93  out << ((i != 31) ? ", " : "");
94  }
95  return out;
96 }
97 
99 template <>
100 inline std::ostream &operator<<(
101  std::ostream &out,
102  ScalarIO<cutlass::Vector<cutlass::int4_t, 8> > const &scalar) {
103 
104  for (int i = 0; i < 8; i++) {
105  out << int(scalar.value[i]);
106  out << ((i != 7) ? ", " : "");
107  }
108  return out;
109 }
110 
112 template <>
113 inline std::ostream &operator<<(
114  std::ostream &out,
115  ScalarIO<cutlass::Vector<cutlass::uint4_t, 8> > const &scalar) {
116 
117  for (int i = 0; i < 8; i++) {
118  out << unsigned(scalar.value[i]);
119  out << ((i != 7) ? ", " : "");
120  }
121  return out;
122 }
123 
125 
126 } // namespace cutlass
cutlass
Definition: convert.h:33
+
coord.h
A Coord is a coordinate of arbitrary rank into a tensor or matrix.
+
cutlass::ScalarIO::ScalarIO
ScalarIO(T value)
Constructs from a value.
Definition: core_io.h:62
+
cutlass::ScalarIO::ScalarIO
ScalarIO()
Default ctor.
Definition: core_io.h:59
+
cutlass::operator<<
std::ostream & operator<<(std::ostream &out, Coord< Rank > const &coord)
Definition: core_io.h:42
+
cutlass::ScalarIO
Helper to enable formatted printing of CUTLASS scalar types to an ostream.
Definition: core_io.h:53
+
cutlass::Vector
Definition: vector.h:62
+
cutlass::ScalarIO::value
T value
Value to print.
Definition: core_io.h:56
+
vector.h
Defines a 1D vector of elements held in the registers of each thread.
-Generated on Thu May 3 2018 16:36:12 for Cutlass by   +Generated on Wed Sep 19 2018 13:58:50 for Cutlass by   doxygen 1.8.14
diff --git a/docs/cutlass_8h.html b/docs/cutlass_8h.html index bbb0463c..419c9123 100644 --- a/docs/cutlass_8h.html +++ b/docs/cutlass_8h.html @@ -73,8 +73,10 @@ $(function() {
+Classes | Namespaces | -Macros
+Macros | +Functions
cutlass.h File Reference
@@ -85,6 +87,13 @@ $(function() {

Go to the source code of this file.

+ + + + + +

+Classes

struct  DebugType< T >
 
struct  DebugValue< Value >
 
@@ -96,18 +105,26 @@ Macros - + + + - - + + +

Namespaces

 cutlass
 
#define CUTLASS_MINOR   0
 
#define CUTLASS_PATCH   0
#define CUTLASS_PATCH   1
 
#define CUTLASS_VERSION   ((CUTLASS_MAJOR)*100 + (CUTLASS_MINOR)*10 + CUTLASS_PATCH)
 
#define CUTLASS_HOST_DEVICE
 
#define CUTLASS_ASSERT(x)   assert(x)
 
#define CUTLASS_PRAGMA_UNROLL
 
#define CUTLASS_PRAGMA_NO_UNROLL
 
#define CUTLASS_ASSERT(x)   assert(x)
 
#define CUTLASS_GEMM_LOOP   CUTLASS_PRAGMA_NO_UNROLL
 
+ + + +

+Functions

template<typename T >
void DebugTypeFunc (T const &t)
 

Macro Definition Documentation

@@ -126,6 +143,20 @@ Macros
+
+
+ +

◆ CUTLASS_GEMM_LOOP

+ +
+
+ + + + +
#define CUTLASS_GEMM_LOOP   CUTLASS_PRAGMA_NO_UNROLL
+
+
@@ -177,7 +208,7 @@ Macros
- +
#define CUTLASS_PATCH   0#define CUTLASS_PATCH   1
@@ -224,12 +255,33 @@ Macros
+
+
+

Function Documentation

+ +

◆ DebugTypeFunc()

+ +
+
+
+template<typename T >
+ + + + + + + + +
void DebugTypeFunc (T const & t)
+
+
-Generated on Thu May 3 2018 16:36:12 for Cutlass by   +Generated on Wed Sep 19 2018 13:58:50 for Cutlass by   doxygen 1.8.14
diff --git a/docs/cutlass_8h_source.html b/docs/cutlass_8h_source.html index d2f44229..9c9fb2b2 100644 --- a/docs/cutlass_8h_source.html +++ b/docs/cutlass_8h_source.html @@ -76,11 +76,14 @@ $(function() {
cutlass.h
-Go to the documentation of this file.
1 /***************************************************************************************************
2  * Copyright (c) 2017-2018, NVIDIA CORPORATION. All rights reserved.
3  *
4  * Redistribution and use in source and binary forms, with or without modification, are permitted
5  * provided that the following conditions are met:
6  * * Redistributions of source code must retain the above copyright notice, this list of
7  * conditions and the following disclaimer.
8  * * Redistributions in binary form must reproduce the above copyright notice, this list of
9  * conditions and the following disclaimer in the documentation and/or other materials
10  * provided with the distribution.
11  * * Neither the name of the NVIDIA CORPORATION nor the names of its contributors may be used
12  * to endorse or promote products derived from this software without specific prior written
13  * permission.
14  *
15  * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND ANY EXPRESS OR
16  * IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND
17  * FITNESS FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE
18  * FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING,
19  * BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS;
20  * OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT,
21  * STRICT LIABILITY, OR TOR (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
22  * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
23  *
24  **************************************************************************************************/
25 
30 #pragma once
31 
33 
34 #define CUTLASS_MAJOR 1
35 #define CUTLASS_MINOR 0
36 #define CUTLASS_PATCH 0
37 #define CUTLASS_VERSION ((CUTLASS_MAJOR)*100 + (CUTLASS_MINOR)*10 + CUTLASS_PATCH)
38 
39 #ifdef __NVCC__
40 #define CUTLASS_HOST_DEVICE __forceinline__ __device__ __host__
41 #define CUTLASS_DEVICE __forceinline__ __device__
42 #elif defined(__CUDACC_RTC__)
43 #define CUTLASS_HOST_DEVICE __forceinline__ __device__
44 #define CUTLASS_DEVICE __forceinline__ __device__
45 #else
46 #define CUTLASS_HOST_DEVICE
47 // CUTLASS_DEVICE is an error if not compiling device code
48 #endif
49 
50 // CUTLASS_PRAGMA_UNROLL inserts a CUTLASS_PRAGMA_UNROLL if supported by the compiler
51 #if defined(__CUDA_ARCH__)
52 #if defined(_MSC_VER)
53 #define CUTLASS_PRAGMA_UNROLL __pragma("unroll")
54 #define CUTLASS_PRAGMA_NO_UNROLL __pragma("unroll 1")
55 #else
56 #define CUTLASS_PRAGMA_UNROLL _Pragma("unroll")
57 #define CUTLASS_PRAGMA_NO_UNROLL _Pragma("unroll 1")
58 #endif
59 #else
60 #define CUTLASS_PRAGMA_UNROLL
61 #define CUTLASS_PRAGMA_NO_UNROLL
62 #endif
63 
64 #define CUTLASS_ASSERT(x) assert(x)
65 
66 namespace cutlass {
67 
69 static const int kWarpSize = 32;
70 
71 } // namespace cutlass
72 
cutlass
Definition: convert.h:33
+Go to the documentation of this file.
1 /***************************************************************************************************
2  * Copyright (c) 2017-2018, NVIDIA CORPORATION. All rights reserved.
3  *
4  * Redistribution and use in source and binary forms, with or without modification, are permitted
5  * provided that the following conditions are met:
6  * * Redistributions of source code must retain the above copyright notice, this list of
7  * conditions and the following disclaimer.
8  * * Redistributions in binary form must reproduce the above copyright notice, this list of
9  * conditions and the following disclaimer in the documentation and/or other materials
10  * provided with the distribution.
11  * * Neither the name of the NVIDIA CORPORATION nor the names of its contributors may be used
12  * to endorse or promote products derived from this software without specific prior written
13  * permission.
14  *
15  * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND ANY EXPRESS OR
16  * IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND
17  * FITNESS FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE
18  * FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING,
19  * BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS;
20  * OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT,
21  * STRICT LIABILITY, OR TOR (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
22  * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
23  *
24  **************************************************************************************************/
25 
30 #pragma once
31 
33 
34 #define CUTLASS_MAJOR 1
35 #define CUTLASS_MINOR 0
36 #define CUTLASS_PATCH 1
37 #define CUTLASS_VERSION ((CUTLASS_MAJOR)*100 + (CUTLASS_MINOR)*10 + CUTLASS_PATCH)
38 
39 #ifdef __NVCC__
40 #define CUTLASS_HOST_DEVICE __forceinline__ __device__ __host__
41 #define CUTLASS_DEVICE __forceinline__ __device__
42 #elif defined(__CUDACC_RTC__)
43 #define CUTLASS_HOST_DEVICE __forceinline__ __device__
44 #define CUTLASS_DEVICE __forceinline__ __device__
45 #else
46 #define CUTLASS_HOST_DEVICE
47 // CUTLASS_DEVICE is an error if not compiling device code
48 #endif
49 
50 #define CUTLASS_ASSERT(x) assert(x)
51 
52 // CUTLASS_PRAGMA_(UNROLL|NO_UNROLL) optimization directives for the CUDA compiler.
53 #if defined(__CUDA_ARCH__)
54 #if defined(_MSC_VER)
55 #define CUTLASS_PRAGMA_UNROLL __pragma("unroll")
56 #define CUTLASS_PRAGMA_NO_UNROLL __pragma("unroll 1")
57 #else
58 #define CUTLASS_PRAGMA_UNROLL _Pragma("unroll")
59 #define CUTLASS_PRAGMA_NO_UNROLL _Pragma("unroll 1")
60 #endif
61 #else
62 #define CUTLASS_PRAGMA_UNROLL
63 #define CUTLASS_PRAGMA_NO_UNROLL
64 #endif
65 
66 #define CUTLASS_GEMM_LOOP CUTLASS_PRAGMA_NO_UNROLL
67 
68 // A small helper class to dump a type at compile time
69 // Usage:: DumpType<Class>::Class
70 template <typename T>
71 struct DebugType {};
72 
73 template <typename T>
74 void DebugTypeFunc(T const& t) {
75  T::t;
76 }
77 
78 // A small helper class to dump a compile time constant at compile time
79 // Usage: DumpValue<Class::kConstant>::kConstant
80 template <int Value>
81 struct DebugValue {};
82 
83 namespace cutlass {
84 
86 static const int kWarpSize = 32;
87 
88 } // namespace cutlass
89 
cutlass
Definition: convert.h:33
+
DebugValue
Definition: cutlass.h:81
+
DebugType
Definition: cutlass.h:71
+
DebugTypeFunc
void DebugTypeFunc(T const &t)
Definition: cutlass.h:74
-Generated on Thu May 3 2018 16:36:12 for Cutlass by   +Generated on Wed Sep 19 2018 13:58:50 for Cutlass by   doxygen 1.8.14
diff --git a/docs/cutlass__math_8h.html b/docs/cutlass__math_8h.html index 953b0d4c..c4dbc54b 100644 --- a/docs/cutlass__math_8h.html +++ b/docs/cutlass__math_8h.html @@ -83,7 +83,7 @@ $(function() {

Math utilities. More...

-
#include <cutlass/util/platform.h>
+
#include "cutlass/util/platform.h"

Go to the source code of this file.

@@ -103,6 +103,10 @@ Classes + + + +
 
struct  cutlass::divide_assert< Dividend, Divisor >
 
struct  cutlass::Min< A, B >
 
struct  cutlass::Max< A, B >
 
@@ -120,11 +124,17 @@ Functions + + + + + +

Namespaces

template<typename value_t >
CUTLASS_HOST_DEVICE value_t cutlass::lcm (value_t a, value_t b)
 
template<typename value_t >
CUTLASS_HOST_DEVICE value_t cutlass::clz (value_t x)
 
template<typename value_t >
CUTLASS_HOST_DEVICE value_t cutlass::find_log2 (value_t x)
 
-Generated on Thu May 3 2018 16:36:12 for Cutlass by   +Generated on Wed Sep 19 2018 13:58:50 for Cutlass by   doxygen 1.8.14
diff --git a/docs/cutlass__math_8h_source.html b/docs/cutlass__math_8h_source.html index 2809a845..8381f641 100644 --- a/docs/cutlass__math_8h_source.html +++ b/docs/cutlass__math_8h_source.html @@ -76,27 +76,33 @@ $(function() {
cutlass_math.h
-Go to the documentation of this file.
1 /***************************************************************************************************
2  * Copyright (c) 2017-2018, NVIDIA CORPORATION. All rights reserved.
3  *
4  * Redistribution and use in source and binary forms, with or without modification, are permitted
5  * provided that the following conditions are met:
6  * * Redistributions of source code must retain the above copyright notice, this list of
7  * conditions and the following disclaimer.
8  * * Redistributions in binary form must reproduce the above copyright notice, this list of
9  * conditions and the following disclaimer in the documentation and/or other materials
10  * provided with the distribution.
11  * * Neither the name of the NVIDIA CORPORATION nor the names of its contributors may be used
12  * to endorse or promote products derived from this software without specific prior written
13  * permission.
14  *
15  * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND ANY EXPRESS OR
16  * IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND
17  * FITNESS FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE
18  * FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING,
19  * BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS;
20  * OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT,
21  * STRICT LIABILITY, OR TOR (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
22  * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
23  *
24  **************************************************************************************************/
25 
26 #pragma once
27 
33 #include <cutlass/util/platform.h>
34 
35 namespace cutlass {
36 
37 /******************************************************************************
38  * Static math utilities
39  ******************************************************************************/
40 
44 template <int N>
45 struct is_pow2 : platform::integral_constant<bool, (N & (N - 1)) == 0> {};
46 
50 template <int N, int CurrentVal = N, int Count = 0>
51 struct log2_down {
53  enum { value = log2_down<N, (CurrentVal >> 1), Count + 1>::value };
54 };
55 
56 // Base case
57 template <int N, int Count>
58 struct log2_down<N, 1, Count> {
59  enum { value = Count };
60 };
61 
65 template <int N, int CurrentVal = N, int Count = 0>
66 struct log2_up {
68  enum { value = log2_up<N, (CurrentVal >> 1), Count + 1>::value };
69 };
70 
71 // Base case
72 template <int N, int Count>
73 struct log2_up<N, 1, Count> {
74  enum { value = ((1 << Count) < N) ? Count + 1 : Count };
75 };
76 
80 template <int N>
81 struct sqrt_est {
82  enum { value = 1 << (log2_up<N>::value / 2) };
83 };
84 
89 template <int Dividend, int Divisor>
90 struct divide_assert {
91  enum { value = Dividend / Divisor };
92 
93  static_assert((Dividend % Divisor == 0), "Not an even multiple");
94 };
95 
96 /******************************************************************************
97  * Rounding
98  ******************************************************************************/
99 
103 template <typename dividend_t, typename divisor_t>
104 CUTLASS_HOST_DEVICE dividend_t round_nearest(dividend_t dividend, divisor_t divisor) {
105  return ((dividend + divisor - 1) / divisor) * divisor;
106 }
107 
111 template <typename value_t>
112 CUTLASS_HOST_DEVICE value_t gcd(value_t a, value_t b) {
113  for (;;) {
114  if (a == 0) return b;
115  b %= a;
116  if (b == 0) return a;
117  a %= b;
118  }
119 }
120 
124 template <typename value_t>
125 CUTLASS_HOST_DEVICE value_t lcm(value_t a, value_t b) {
126  value_t temp = gcd(a, b);
127 
128  return temp ? (a / temp * b) : 0;
129 }
130 
131 } // namespace cutlass
cutlass::divide_assert::value
Definition: cutlass_math.h:91
+Go to the documentation of this file.
1 /***************************************************************************************************
2  * Copyright (c) 2017-2018, NVIDIA CORPORATION. All rights reserved.
3  *
4  * Redistribution and use in source and binary forms, with or without modification, are permitted
5  * provided that the following conditions are met:
6  * * Redistributions of source code must retain the above copyright notice, this list of
7  * conditions and the following disclaimer.
8  * * Redistributions in binary form must reproduce the above copyright notice, this list of
9  * conditions and the following disclaimer in the documentation and/or other materials
10  * provided with the distribution.
11  * * Neither the name of the NVIDIA CORPORATION nor the names of its contributors may be used
12  * to endorse or promote products derived from this software without specific prior written
13  * permission.
14  *
15  * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND ANY EXPRESS OR
16  * IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND
17  * FITNESS FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE
18  * FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING,
19  * BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS;
20  * OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT,
21  * STRICT LIABILITY, OR TOR (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
22  * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
23  *
24  **************************************************************************************************/
25 
26 #pragma once
27 
33 #include "cutlass/util/platform.h"
34 
35 namespace cutlass {
36 
37 /******************************************************************************
38  * Static math utilities
39  ******************************************************************************/
40 
44 template <int N>
45 struct is_pow2 : platform::integral_constant<bool, (N & (N - 1)) == 0> {};
46 
50 template <int N, int CurrentVal = N, int Count = 0>
51 struct log2_down {
53  enum { value = log2_down<N, (CurrentVal >> 1), Count + 1>::value };
54 };
55 
56 // Base case
57 template <int N, int Count>
58 struct log2_down<N, 1, Count> {
59  enum { value = Count };
60 };
61 
65 template <int N, int CurrentVal = N, int Count = 0>
66 struct log2_up {
68  enum { value = log2_up<N, (CurrentVal >> 1), Count + 1>::value };
69 };
70 
71 // Base case
72 template <int N, int Count>
73 struct log2_up<N, 1, Count> {
74  enum { value = ((1 << Count) < N) ? Count + 1 : Count };
75 };
76 
80 template <int N>
81 struct sqrt_est {
82  enum { value = 1 << (log2_up<N>::value / 2) };
83 };
84 
89 template <int Dividend, int Divisor>
90 struct divide_assert {
91  enum { value = Dividend / Divisor };
92 
93  static_assert((Dividend % Divisor == 0), "Not an even multiple");
94 };
95 
96 /******************************************************************************
97  * Rounding
98  ******************************************************************************/
99 
103 template <typename dividend_t, typename divisor_t>
104 CUTLASS_HOST_DEVICE dividend_t round_nearest(dividend_t dividend, divisor_t divisor) {
105  return ((dividend + divisor - 1) / divisor) * divisor;
106 }
107 
111 template <typename value_t>
112 CUTLASS_HOST_DEVICE value_t gcd(value_t a, value_t b) {
113  for (;;) {
114  if (a == 0) return b;
115  b %= a;
116  if (b == 0) return a;
117  a %= b;
118  }
119 }
120 
124 template <typename value_t>
125 CUTLASS_HOST_DEVICE value_t lcm(value_t a, value_t b) {
126  value_t temp = gcd(a, b);
127 
128  return temp ? (a / temp * b) : 0;
129 }
130 
136 template <typename value_t>
137 CUTLASS_HOST_DEVICE value_t clz(value_t x) {
138  for (int i = 31; i >= 0; --i) {
139  if ((1 << i) & x) return 31 - i;
140  }
141  return 32;
142 }
143 
144 template <typename value_t>
145 CUTLASS_HOST_DEVICE value_t find_log2(value_t x) {
146  int a = 31 - clz(x);
147  a += (x & (x - 1)) != 0; // Round up, add 1 if not a power of 2.
148  return a;
149 }
150 
151 /******************************************************************************
152  * Min/Max
153  ******************************************************************************/
154 
155 template <int A, int B>
156 struct Min {
157  static int const kValue = (A < B) ? A : B;
158 };
159 
160 template <int A, int B>
161 struct Max {
162  static int const kValue = (A > B) ? A : B;
163 };
164 
165 } // namespace cutlass
cutlass::divide_assert::value
Definition: cutlass_math.h:91
cutlass
Definition: convert.h:33
+
cutlass::Min::kValue
static int const kValue
Definition: cutlass_math.h:157
+
cutlass::find_log2
CUTLASS_HOST_DEVICE value_t find_log2(value_t x)
Definition: cutlass_math.h:145
cutlass::log2_down
Definition: cutlass_math.h:51
platform.h
C++ features that may be otherwise unimplemented for CUDA device functions.
+
cutlass::Min
Definition: cutlass_math.h:156
cutlass::log2_down::value
Definition: cutlass_math.h:53
cutlass::lcm
CUTLASS_HOST_DEVICE value_t lcm(value_t a, value_t b)
Definition: cutlass_math.h:125
cutlass::round_nearest
CUTLASS_HOST_DEVICE dividend_t round_nearest(dividend_t dividend, divisor_t divisor)
Definition: cutlass_math.h:104
cutlass::log2_up::value
Definition: cutlass_math.h:68
-
cutlass::platform::integral_constant
std::integral_constant
Definition: platform.h:274
+
cutlass::platform::integral_constant
std::integral_constant
Definition: platform.h:282
CUTLASS_HOST_DEVICE
#define CUTLASS_HOST_DEVICE
Definition: cutlass.h:46
-
static_assert
#define static_assert(__e, __m)
Definition: platform.h:145
+
static_assert
#define static_assert(__e, __m)
Definition: platform.h:153
+
cutlass::Max
Definition: cutlass_math.h:161
cutlass::sqrt_est::value
Definition: cutlass_math.h:82
cutlass::gcd
CUTLASS_HOST_DEVICE value_t gcd(value_t a, value_t b)
Definition: cutlass_math.h:112
cutlass::divide_assert
Definition: cutlass_math.h:90
cutlass::log2_up
Definition: cutlass_math.h:66
+
cutlass::clz
CUTLASS_HOST_DEVICE value_t clz(value_t x)
Definition: cutlass_math.h:137
cutlass::is_pow2
Definition: cutlass_math.h:45
+
cutlass::Max::kValue
static int const kValue
Definition: cutlass_math.h:162
cutlass::sqrt_est
Definition: cutlass_math.h:81
-Generated on Thu May 3 2018 16:36:12 for Cutlass by   +Generated on Wed Sep 19 2018 13:58:50 for Cutlass by   doxygen 1.8.14
diff --git a/docs/debug_8h.html b/docs/debug_8h.html index 1f88396a..81ed9f3c 100644 --- a/docs/debug_8h.html +++ b/docs/debug_8h.html @@ -231,7 +231,7 @@ Functions
- + - +

Classes

struct  cutlass::gemm::DgemmConfig< OutputTile_, AccumulatorsPerThread_, kScalarsPerLdgA_, kScalarsPerLdgB_ >
struct  cutlass::gemm::DgemmConfig< OutputTile_, ThreadGemmShape_, kScalarsPerLdgA_, kScalarsPerLdgB_ >
 
struct  cutlass::gemm::DgemmTraits< kLayoutA_, kLayoutB_, OutputTile_, EpilogueFunctor_, AccumulatorsPerThread_, kScalarsPerLdgA_, kScalarsPerLdgB_, Index_, GemmConfig_, GemmEpilogueTraits_ >
struct  cutlass::gemm::DgemmTraits< kLayoutA_, kLayoutB_, OutputTile_, EpilogueFunctor_, ThreadGemmShape_, kScalarsPerLdgA_, kScalarsPerLdgB_, Index_, GemmConfig_, GemmEpilogueTraits_ >
 
-Generated on Thu May 3 2018 16:36:12 for Cutlass by   +Generated on Wed Sep 19 2018 13:58:50 for Cutlass by   doxygen 1.8.14
diff --git a/docs/dgemm__traits_8h_source.html b/docs/dgemm__traits_8h_source.html index 9cf2c873..d7cdbe52 100644 --- a/docs/dgemm__traits_8h_source.html +++ b/docs/dgemm__traits_8h_source.html @@ -76,26 +76,26 @@ $(function() {
dgemm_traits.h
-Go to the documentation of this file.
1 /***************************************************************************************************
2  * Copyright (c) 2017-2018, NVIDIA CORPORATION. All rights reserved.
3  *
4  * Redistribution and use in source and binary forms, with or without modification, are permitted
5  * provided that the following conditions are met:
6  * * Redistributions of source code must retain the above copyright notice, this list of
7  * conditions and the following disclaimer.
8  * * Redistributions in binary form must reproduce the above copyright notice, this list of
9  * conditions and the following disclaimer in the documentation and/or other materials
10  * provided with the distribution.
11  * * Neither the name of the NVIDIA CORPORATION nor the names of its contributors may be used
12  * to endorse or promote products derived from this software without specific prior written
13  * permission.
14  *
15  * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND ANY EXPRESS OR
16  * IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND
17  * FITNESS FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE
18  * FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING,
19  * BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS;
20  * OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT,
21  * STRICT LIABILITY, OR TOR (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
22  * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
23  *
24  **************************************************************************************************/
28 #pragma once
29 
30 #include <cutlass/gemm/gemm.h>
31 #include <cutlass/gemm/gemm_epilogue.h>
32 #include <cutlass/gemm/gemm_epilogue_traits.h>
33 #include <cutlass/gemm/gemm_global_tile.h>
34 #include <cutlass/gemm/gemm_shared_tile.h>
35 #include <cutlass/gemm/gemm_traits.h>
36 #include <cutlass/gemm/thread_multiply_add.h>
37 
38 namespace cutlass {
39 namespace gemm {
40 
42 
43 template <
45  typename OutputTile_,
47  typename AccumulatorsPerThread_,
49  int kScalarsPerLdgA_ = 1,
51  int kScalarsPerLdgB_ = 1>
52 struct DgemmConfig
53  : public GemmConfig<
55  double,
57  double,
59  double,
61  double,
63  OutputTile_,
65  ThreadMultiplyAdd<AccumulatorsPerThread_, Shape<1, 4, 8>, double, double, double>,
67  kScalarsPerLdgA_,
69  kScalarsPerLdgA_,
71  2,
73  kScalarsPerLdgB_,
75  kScalarsPerLdgB_,
77  2,
79  1,
81  2,
83  1,
85  2> {};
86 
88 
89 template <
91  MatrixLayout::Kind kLayoutA_,
93  MatrixLayout::Kind kLayoutB_,
95  typename OutputTile_ = Shape<8, 64, 128>,
97  typename EpilogueFunctor_ = LinearScaling<double>,
99  typename AccumulatorsPerThread_ = Shape<8, 8, 8>,
101  int kScalarsPerLdgA_ = 1,
103  int kScalarsPerLdgB_ = 1,
105  typename Index_ = int,
107  typename GemmConfig_ =
108  DgemmConfig<OutputTile_, AccumulatorsPerThread_, kScalarsPerLdgA_, kScalarsPerLdgB_>,
110  typename GemmEpilogueTraits_ =
111  SimplifiedGemmEpilogueTraits<GemmConfig_, EpilogueFunctor_, Index_> >
112 struct DgemmTraits : public SimplifiedGemmTraits<
113  // The layout for A.
114  kLayoutA_,
115  // The layout for B.
116  kLayoutB_,
117  // The config.
118  GemmConfig_,
119  // The epilogue.
120  GemmEpilogue<GemmEpilogueTraits_>,
121  // The index.
122  Index_> {};
123 
125 
126 } // namespace gemm
127 } // namespace cutlass
cutlass
Definition: convert.h:33
+Go to the documentation of this file.
1 /***************************************************************************************************
2  * Copyright (c) 2017-2018, NVIDIA CORPORATION. All rights reserved.
3  *
4  * Redistribution and use in source and binary forms, with or without modification, are permitted
5  * provided that the following conditions are met:
6  * * Redistributions of source code must retain the above copyright notice, this list of
7  * conditions and the following disclaimer.
8  * * Redistributions in binary form must reproduce the above copyright notice, this list of
9  * conditions and the following disclaimer in the documentation and/or other materials
10  * provided with the distribution.
11  * * Neither the name of the NVIDIA CORPORATION nor the names of its contributors may be used
12  * to endorse or promote products derived from this software without specific prior written
13  * permission.
14  *
15  * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND ANY EXPRESS OR
16  * IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND
17  * FITNESS FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE
18  * FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING,
19  * BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS;
20  * OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT,
21  * STRICT LIABILITY, OR TOR (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
22  * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
23  *
24  **************************************************************************************************/
28 #pragma once
29 
30 #include "cutlass/gemm/gemm.h"
31 #include "cutlass/gemm/gemm_epilogue.h"
32 #include "cutlass/gemm/gemm_epilogue_traits.h"
33 #include "cutlass/gemm/gemm_global_tile.h"
34 #include "cutlass/gemm/gemm_shared_tile.h"
35 #include "cutlass/gemm/gemm_traits.h"
36 #include "cutlass/gemm/thread_multiply_add.h"
37 
38 namespace cutlass {
39 namespace gemm {
40 
42 
43 template <
45  typename OutputTile_,
47  typename ThreadGemmShape_,
49  int kScalarsPerLdgA_ = 1,
51  int kScalarsPerLdgB_ = 1>
52 struct DgemmConfig
53  : public GemmConfig<
55  double,
57  double,
59  double,
61  double,
63  OutputTile_,
65  ThreadMultiplyAdd<ThreadGemmShape_, Shape<1, 4, 8>, double, double, double>,
67  kScalarsPerLdgA_,
69  kScalarsPerLdgA_,
71  2,
73  kScalarsPerLdgB_,
75  kScalarsPerLdgB_,
77  2,
79  1,
81  2,
83  1,
85  2,
87  false,
89  false,
91  false
92  >{};
93 
95 
96 template <
98  MatrixLayout::Kind kLayoutA_,
100  MatrixLayout::Kind kLayoutB_,
102  typename OutputTile_ = Shape<8, 64, 128>,
104  typename EpilogueFunctor_ = LinearScaling<double>,
106  typename ThreadGemmShape_ = Shape<8, 8, 8>,
108  int kScalarsPerLdgA_ = 1,
110  int kScalarsPerLdgB_ = 1,
112  typename Index_ = int,
114  typename GemmConfig_ =
115  DgemmConfig<OutputTile_, ThreadGemmShape_, kScalarsPerLdgA_, kScalarsPerLdgB_>,
117  typename GemmEpilogueTraits_ =
118  SimplifiedGemmEpilogueTraits<GemmConfig_, EpilogueFunctor_, Index_> >
119 struct DgemmTraits : public SimplifiedGemmTraits<
120  // The layout for A.
121  kLayoutA_,
122  // The layout for B.
123  kLayoutB_,
124  // The config.
125  GemmConfig_,
126  // The epilogue.
127  GemmEpilogue<GemmEpilogueTraits_>,
128  // The index.
129  Index_> {};
130 
132 
133 } // namespace gemm
134 } // namespace cutlass
cutlass
Definition: convert.h:33
gemm_global_tile.h
Defines iterators for efficiently loading and storing to global memory.
gemm_traits.h
Defines structural properties of complete GEMM computation.
thread_multiply_add.h
Template implementing matrix multiply-add operations on fragments.
gemm_epilogue.h
Implements the epilogue phase of the GEMM kernel that efficiently updates global memory with the comp...
gemm_shared_tile.h
Defines iterators for efficiently loading and storing tiles to and from shared memory.
-
cutlass::gemm::GemmConfig
Definition: gemm_traits.h:79
-
cutlass::gemm::DgemmTraits
Definition: dgemm_traits.h:112
+
cutlass::gemm::GemmConfig
Definition: gemm_config.h:76
+
cutlass::gemm::DgemmTraits
Definition: dgemm_traits.h:119
cutlass::gemm::DgemmConfig
Definition: dgemm_traits.h:52
cutlass::Shape
A Shape implementing Layout Concept describing the dimensions of a cube.
Definition: shape.h:64
-
cutlass::gemm::SimplifiedGemmEpilogueTraits
Definition: gemm_epilogue_traits.h:300
-
cutlass::MatrixLayout::Kind
Kind
Definition: matrix_traits.h:36
-
cutlass::gemm::LinearScaling
Functor to compute linear combination of fragments.
Definition: linear_scaling.h:40
+
cutlass::gemm::SimplifiedGemmEpilogueTraits
Definition: gemm_epilogue_traits.h:323
+
cutlass::MatrixLayout::Kind
Kind
Enumeration defining fundamental contiguous layouts.
Definition: matrix_traits.h:159
+
cutlass::gemm::LinearScaling
Functor to compute linear combination of fragments.
Definition: linear_scaling.h:51
gemm.h
Implements a software-pipelined efficient GEMM.
gemm_epilogue_traits.h
Defines structural properties of the GEMM epilogue.
-
cutlass::gemm::SimplifiedGemmTraits
Definition: gemm_traits.h:723
+
cutlass::gemm::SimplifiedGemmTraits
Definition: gemm_traits.h:650
-Generated on Thu May 3 2018 16:36:12 for Cutlass by   +Generated on Wed Sep 19 2018 13:58:50 for Cutlass by   doxygen 1.8.14
diff --git a/docs/dir_1417ee5ebebc309c36b7962f26a92c39.html b/docs/dir_1417ee5ebebc309c36b7962f26a92c39.html index d7393ef1..6555e36c 100644 --- a/docs/dir_1417ee5ebebc309c36b7962f26a92c39.html +++ b/docs/dir_1417ee5ebebc309c36b7962f26a92c39.html @@ -101,15 +101,15 @@ Files
- - - + + + @@ -128,12 +128,24 @@ Files + + + + + + + + + + + + @@ -143,11 +155,20 @@ Files + + + + + + + + +

@@ -109,7 +109,7 @@ Namespaces

file  fragment.h [code]
 Defines Fragment, a statically-sized array for storing parts of matrices within a thread's registers.
 
file  fragment_load_store.h [code]
 Defines accessors for loading and storing fragments to memory efficiently.
 
file  fragment_multiply_add.h [code]
 Defines multiply-add operations on fragments within a thread.
 
file  iterator_access.h [code]
 Free functions for loading and storing to implementations of tile iteartor concepts.
 
file  kernel_launch.h [code]
 Defines structures and helpers to launch CUDA kernels within CUTLASS.
 
file  load_store.h [code]
 Defines abstractions for efficiently loading and storing vectors to memory.
 
file  tensor_ref.h [code]
 Defines a structure containing strides, bounds, and a pointer to tensor data.
 
file  tensor_ref_collection.h [code]
 Introduces TensorRefCollection concept and defines TensorRefBatch and TensorRefArray.
 
file  tensor_view.h [code]
 Defines a structure containing strides and a pointer to tensor data.
 
file  tile_allocation.h [code]
 Defines a fragment based on a Shape<> template.
 
file  tile_coord.h [code]
 Defines a coordinate used for the CUTLASS 4-D tile structure.
 
file  tile_iterator.h [code]
 Defines the Tile Traits concept and iterators for loading and storing to tiles efficiently.
 
file  tile_stream.h [code]
 Implements the tile stream concept, composing an iterator with a transformation. Offers split-phase semantics, separating the initiation of an asynchronous memory operation with a fence forcing it to complete.
 
file  tile_traits_standard.h [code]
 Defines tile traits for several tile partitioning arrangements of threads expected to achieve efficient streaming performance.
 
file  wmma_matrix.h [code]
 Abstractions for loading and storing matrices using the CUDA WMMA API.
 
file  zip_fragment.h [code]
 Models a pair of fragments.
 
file  zip_tensor_ref.h [code]
 Defines a structure containing a pair of TensorRef-like objects.
 
file  zip_tile_iterator.h [code]
 Constructs an iterator that owns two tile iterator instances.
 
-Generated on Thu May 3 2018 16:36:15 for Cutlass by   +Generated on Wed Sep 19 2018 13:58:53 for Cutlass by   doxygen 1.8.14
diff --git a/docs/dir_18d6a367a3982a494d65599933fc67a3.html b/docs/dir_18d6a367a3982a494d65599933fc67a3.html index 16126747..b606ad3e 100644 --- a/docs/dir_18d6a367a3982a494d65599933fc67a3.html +++ b/docs/dir_18d6a367a3982a494d65599933fc67a3.html @@ -85,9 +85,24 @@ Files
file  dgemm_traits.h [code]
 Defines structural traits of double-precision GEMM.
 
file  fp16_sgemm_multiply_add.h [code]
 Template implementing matrix multiply-add operations on fragments.
 
file  fp16_sgemm_traits.h [code]
 Defies structural properties of single-precision GEMM where any number of the input/output could be fp16 or fp32. The accumulator type stays in fp32.
 
file  gemm.h [code]
 Implements a software-pipelined efficient GEMM.
 
file  gemm_config.h [code]
 Defines properties of GEMM computation that impose some constraints on caller.
 
file  gemm_coord.h [code]
 GemmCoord is a structure derived from Coord<4> that specifies a location within the coordinate system of a GEMM problem.
 
file  gemm_desc.h [code]
 Implements a software-pipelined efficient GEMM.
 
file  gemm_epilogue.h [code]
 Implements the epilogue phase of the GEMM kernel that efficiently updates global memory with the computed matrix product.
 
file  gemm_shared_tile.h [code]
 Defines iterators for efficiently loading and storing tiles to and from shared memory.
 
file  gemm_stream_pair.h [code]
 Defines a pair of GEMM tile streams.
 
file  gemm_traits.h [code]
 Defines structural properties of complete GEMM computation.
 
file  hgemm_traits.h [code]
 Defies structural properties of half-precision GEMM computation.
 
file  identity_block_swizzle.h [code]
 Defies functors for mapping blockIdx to partitions of the GEMM computation.
 
file  igemm_epilogue.h [code]
 Defines the epilogue phase of the GEMM computation for IGEMM, supporting integer and floating-point output matrix formats.
 
file  linear_scaling.h [code]
 Implements the BLAS linear scaling function alpha*AB + beta*C.
 
file  linear_scaling_device_ptr.h [code]
 Implements the BLAS linear scaling function alpha*AB + beta*C.
 
file  scalar_or_pointer.h [code]
 Implements the BLAS linear scaling function alpha*AB + beta*C.
 
file  sgemm_traits.h [code]
 Defies structural properties of single-precision GEMM.
 
file  thread_multiply_add.h [code]
 Template implementing matrix multiply-add operations on fragments.
 
file  threadblock_swizzle.h [code]
 Defies functors for mapping blockIdx to partitions of the GEMM computation.
 
file  wmma_gemm_epilogue_traits.h [code]
 Defines structural properties of WMMA GEMM's epilogue phase.
 
+ + + + @@ -92,7 +96,7 @@ Files
-Generated on Thu May 3 2018 16:36:15 for Cutlass by   +Generated on Wed Sep 19 2018 13:58:53 for Cutlass by   doxygen 1.8.14
diff --git a/docs/files.html b/docs/files.html index 2c06de5a..10195209 100644 --- a/docs/files.html +++ b/docs/files.html @@ -75,62 +75,79 @@ $(function() {
Here is a list of all files with brief descriptions:

Files

file  complex.h [code]
 
file  cutlass_math.h [code]
 Math utilities.
 
file  debug.h [code]
 Debugging and logging functionality.
 
file  numeric_types.h [code]
 
file  platform.h [code]
 C++ features that may be otherwise unimplemented for CUDA device functions.
 
- - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + +
 clear_accumulators.hDefines abstractions for efficiently clearing accumulator tiles
 convert.hDefines conversion operations among Fragments of different base type
 coord.hA Coord is a coordinate of arbitrary rank into a tensor or matrix
 core_io.hHelpers for printing cutlass/core objects
 cutlass.hBasic include for CUTLASS macros
 cutlass_math.hMath utilities
 debug.hDebugging and logging functionality
 dgemm_traits.hDefines structural traits of double-precision GEMM
 fragment.hDefines Fragment, a statically-sized array for storing parts of matrices within a thread's registers
 fragment_load_store.hDefines accessors for loading and storing fragments to memory efficiently
 fragment_multiply_add.hDefines multiply-add operations on fragments within a thread
 gemm.hImplements a software-pipelined efficient GEMM
 gemm_epilogue.hImplements the epilogue phase of the GEMM kernel that efficiently updates global memory with the computed matrix product
 gemm_epilogue_traits.hDefines structural properties of the GEMM epilogue
 gemm_global_stream.hImplements efficient loading of the thread block-level tile from global memory and storing to shared memory
 gemm_global_tile.hDefines iterators for efficiently loading and storing to global memory
 gemm_operand.hDefines constant expressions for mapping GEMM problem size and strides onto pitch-linear memory
 gemm_shared_stream.hDefines abstractions for managing loading and storing fragments to shared memory in the efficient GEMM pipeline
 gemm_shared_tile.hDefines iterators for efficiently loading and storing tiles to and from shared memory
 gemm_traits.hDefines structural properties of complete GEMM computation
 hgemm_global_tile.hTile traits used to construct global tile iterator for HGEMM. This is intended to partition the thread block-level tile into 2D subtiles loaded by the threads and facilitate memory accesses larger than 16 bits
 hgemm_multiply_add.hSpecialization implementing multiply-add operation on half-precision floating point fragments
 hgemm_swizzle.hTransposes a tile of 16b elements. Used by HGEMM to construct a K-strided layout in shared memory for multiplicands
 hgemm_traits.hDefies structural properties of half-precision GEMM computation
 identity_block_swizzle.hDefies functors for mapping blockIdx to partitions of the GEMM computation
 igemm_epilogue.hDefines the epilogue phase of the GEMM computation for IGEMM, supporting integer and floating-point output matrix formats
 igemm_global_tile.hImplements tile iterators to partition the thread block tile into 2D subtiles and efficiently load each. Applies permute transformation to construct 'interleaved K-strided' data layout in which 4-element dot products from the same K index are arranged in consecutive locations within shared memory
 igemm_multiply_add.hImplements matrix multiply accumulate operation of 8-bit integer data using DP4A instruction
 igemm_swizzle.hTransposes a fragment of data containing packed 8-bit integer elements
 igemm_traits.hDefies structural properties of mixed-precision integer GEMM. Multiplicands are assumed to be packed 8bit integers, accumulators are assumed to be 32b signed integers, and output formats vary
 iterator_access.hFree functions for loading and storing to implementations of tile iteartor concepts
 linear_scaling.hImplements the BLAS linear scaling function alpha*AB + beta*C
 load_store.hDefines abstractions for efficiently loading and storing vectors to memory
 matrix_traits.hDefines properties of matrices used to denote layout and operands to GEMM kernels
 platform.hC++ features that may be otherwise unimplemented for CUDA device functions
 predicate_vector.hDefines container classes and iterators for managing a statically sized vector of boolean predicates
 reshape_tile.hDefines a type for restructuring a tile
 sgemm_traits.hDefies structural properties of single-precision GEMM
 shape.hDefines Shape implementing the Layout concept for representing a 4D hypercube of objects
 tensor_ref.hDefines a structure containing strides, bounds, and a pointer to tensor data
 tensor_view.hDefines a structure containing strides and a pointer to tensor data
 thread_multiply_add.hTemplate implementing matrix multiply-add operations on fragments
 tile_iterator.hDefines the Tile Traits concept and iterators for loading and storing to tiles efficiently
 tile_traits_standard.hDefines tile traits for several tile partitioning arrangements of threads expected to achieve efficient streaming performance
 vector.hDefines a 1D vector of elements held in the registers of each thread
 wmma_gemm_epilogue_traits.hDefines structural properties of WMMA GEMM's epilogue phase
 wmma_gemm_global_tile.hDefines tile iterator traits for loading thread block-level tile from global memory
 wmma_gemm_multiply_add.hImplements warp-level matrix multiply-accumulate operation using CUDA WMMA API
 wmma_gemm_shared_tile.hDefines iterator traits for efficiently loading and storing fragment to and from shared memory, specialized for WMMA GEMM
 wmma_gemm_traits.hDefies structural properties of GEMM targeting WMMA API in CUDA
 wmma_matrix.hAbstractions for loading and storing matrices using the CUDA WMMA API
 complex.h
 convert.hDefines conversion operations among Fragments of different base type
 coord.hA Coord is a coordinate of arbitrary rank into a tensor or matrix
 core_io.hHelpers for printing cutlass/core objects
 cutlass.hBasic include for CUTLASS macros
 cutlass_math.hMath utilities
 debug.hDebugging and logging functionality
 dgemm_traits.hDefines structural traits of double-precision GEMM
 fp16_sgemm_multiply_add.hTemplate implementing matrix multiply-add operations on fragments
 fp16_sgemm_traits.hDefies structural properties of single-precision GEMM where any number of the input/output could be fp16 or fp32. The accumulator type stays in fp32
 fragment.hDefines Fragment, a statically-sized array for storing parts of matrices within a thread's registers
 fragment_multiply_add.hDefines multiply-add operations on fragments within a thread
 gemm.hImplements a software-pipelined efficient GEMM
 gemm_config.hDefines properties of GEMM computation that impose some constraints on caller
 gemm_coord.hGemmCoord is a structure derived from Coord<4> that specifies a location within the coordinate system of a GEMM problem
 gemm_desc.hImplements a software-pipelined efficient GEMM
 gemm_epilogue.hImplements the epilogue phase of the GEMM kernel that efficiently updates global memory with the computed matrix product
 gemm_epilogue_traits.hDefines structural properties of the GEMM epilogue
 gemm_global_stream.hImplements efficient loading of the thread block-level tile from global memory and storing to shared memory
 gemm_global_tile.hDefines iterators for efficiently loading and storing to global memory
 gemm_operand.hDefines constant expressions for mapping GEMM problem size and strides onto pitch-linear memory
 gemm_shared_stream.hDefines abstractions for managing loading and storing fragments to shared memory in the efficient GEMM pipeline
 gemm_shared_tile.hDefines iterators for efficiently loading and storing tiles to and from shared memory
 gemm_stream_pair.hDefines a pair of GEMM tile streams
 gemm_traits.hDefines structural properties of complete GEMM computation
 hgemm_global_tile.hTile traits used to construct global tile iterator for HGEMM. This is intended to partition the thread block-level tile into 2D subtiles loaded by the threads and facilitate memory accesses larger than 16 bits
 hgemm_multiply_add.hSpecialization implementing multiply-add operation on half-precision floating point fragments
 hgemm_swizzle.hTransposes a tile of 16b elements. Used by HGEMM to construct a K-strided layout in shared memory for multiplicands
 hgemm_traits.hDefies structural properties of half-precision GEMM computation
 igemm_epilogue.hDefines the epilogue phase of the GEMM computation for IGEMM, supporting integer and floating-point output matrix formats
 igemm_global_tile.hImplements tile iterators to partition the thread block tile into 2D subtiles and efficiently load each. Applies permute transformation to construct 'interleaved K-strided' data layout in which 4-element dot products from the same K index are arranged in consecutive locations within shared memory
 igemm_multiply_add.hImplements matrix multiply accumulate operation of 8-bit integer data using DP4A instruction
 igemm_swizzle.hTransposes a fragment of data containing packed 8-bit integer elements
 igemm_traits.hDefies structural properties of mixed-precision integer GEMM. Multiplicands are assumed to be packed 8bit integers, accumulators are assumed to be 32b signed integers, and output formats vary
 iterator_access.hFree functions for loading and storing to implementations of tile iteartor concepts
 kernel_launch.hDefines structures and helpers to launch CUDA kernels within CUTLASS
 linear_scaling.hImplements the BLAS linear scaling function alpha*AB + beta*C
 linear_scaling_device_ptr.hImplements the BLAS linear scaling function alpha*AB + beta*C
 load_store.hDefines abstractions for efficiently loading and storing vectors to memory
 matrix_traits.hDefines properties of matrices used to denote layout and operands to GEMM kernels
 numeric_types.h
 platform.hC++ features that may be otherwise unimplemented for CUDA device functions
 predicate_vector.hDefines container classes and iterators for managing a statically sized vector of boolean predicates
 reshape_tile.hDefines a type for restructuring a tile
 scalar_or_pointer.hImplements the BLAS linear scaling function alpha*AB + beta*C
 sgemm_traits.hDefies structural properties of single-precision GEMM
 shape.hDefines Shape implementing the Layout concept for representing a 4D hypercube of objects
 tensor_ref.hDefines a structure containing strides, bounds, and a pointer to tensor data
 tensor_ref_collection.hIntroduces TensorRefCollection concept and defines TensorRefBatch and TensorRefArray
 tensor_view.hDefines a structure containing strides and a pointer to tensor data
 thread_multiply_add.hTemplate implementing matrix multiply-add operations on fragments
 threadblock_swizzle.hDefies functors for mapping blockIdx to partitions of the GEMM computation
 tile_allocation.hDefines a fragment based on a Shape<> template
 tile_coord.hDefines a coordinate used for the CUTLASS 4-D tile structure
 tile_iterator.hDefines the Tile Traits concept and iterators for loading and storing to tiles efficiently
 tile_stream.hImplements the tile stream concept, composing an iterator with a transformation. Offers split-phase semantics, separating the initiation of an asynchronous memory operation with a fence forcing it to complete
 tile_traits_standard.hDefines tile traits for several tile partitioning arrangements of threads expected to achieve efficient streaming performance
 vector.hDefines a 1D vector of elements held in the registers of each thread
 wmma_gemm_epilogue_traits.hDefines structural properties of WMMA GEMM's epilogue phase
 wmma_gemm_global_tile.hDefines tile iterator traits for loading thread block-level tile from global memory
 wmma_gemm_multiply_add.hImplements warp-level matrix multiply-accumulate operation using CUDA WMMA API
 wmma_gemm_shared_tile.hDefines iterator traits for efficiently loading and storing fragment to and from shared memory, specialized for WMMA GEMM
 wmma_gemm_traits.hDefies structural properties of GEMM targeting WMMA API in CUDA
 wmma_matrix.hAbstractions for loading and storing matrices using the CUDA WMMA API
 zip_fragment.hModels a pair of fragments
 zip_tensor_ref.hDefines a structure containing a pair of TensorRef-like objects
 zip_tile_iterator.hConstructs an iterator that owns two tile iterator instances
-Generated on Thu May 3 2018 16:36:15 for Cutlass by   +Generated on Wed Sep 19 2018 13:58:53 for Cutlass by   doxygen 1.8.14
diff --git a/docs/fp16__sgemm__multiply__add_8h.html b/docs/fp16__sgemm__multiply__add_8h.html new file mode 100644 index 00000000..deff050e --- /dev/null +++ b/docs/fp16__sgemm__multiply__add_8h.html @@ -0,0 +1,111 @@ + + + + + + + +Cutlass: fp16_sgemm_multiply_add.h File Reference + + + + + + + + + + +
+
+ + + + + + +
+
Cutlass +
+
CUDA Templates for Linear Algebra Subroutines and Solvers
+
+
+ + + + + + + + +
+
+ + +
+ +
+ + +
+
+
+Classes | +Namespaces
+
+
fp16_sgemm_multiply_add.h File Reference
+
+
+ +

Template implementing matrix multiply-add operations on fragments. +More...

+
#include "cutlass/fragment.h"
+#include "cutlass/gemm/thread_multiply_add.h"
+
+

Go to the source code of this file.

+ + + + + +

+Classes

struct  cutlass::gemm::ThreadMultiplyAdd< ThreadGemmShape_, ThreadsPerWarp_, half, half, float >
 Template performing matrix multiply-add operation within a thread. More...
 
+ + + + + +

+Namespaces

 cutlass
 
 cutlass::gemm
 
+
+ +
+Generated on Wed Sep 19 2018 13:58:50 for Cutlass by   +doxygen + 1.8.14 +
+ + diff --git a/docs/fp16__sgemm__multiply__add_8h_source.html b/docs/fp16__sgemm__multiply__add_8h_source.html new file mode 100644 index 00000000..efac0463 --- /dev/null +++ b/docs/fp16__sgemm__multiply__add_8h_source.html @@ -0,0 +1,107 @@ + + + + + + + +Cutlass: fp16_sgemm_multiply_add.h Source File + + + + + + + + + + +
+
+ + + + + + +
+
Cutlass +
+
CUDA Templates for Linear Algebra Subroutines and Solvers
+
+
+ + + + + + + + +
+
+ + +
+ +
+ + +
+
+
+
fp16_sgemm_multiply_add.h
+
+
+Go to the documentation of this file.
1 /***************************************************************************************************
2  * Copyright (c) 2017-2018, NVIDIA CORPORATION. All rights reserved.
3  *
4  * Redistribution and use in source and binary forms, with or without modification, are permitted
5  * provided that the following conditions are met:
6  * * Redistributions of source code must retain the above copyright notice, this list of
7  * conditions and the following disclaimer.
8  * * Redistributions in binary form must reproduce the above copyright notice, this list of
9  * conditions and the following disclaimer in the documentation and/or other materials
10  * provided with the distribution.
11  * * Neither the name of the NVIDIA CORPORATION nor the names of its contributors may be used
12  * to endorse or promote products derived from this software without specific prior written
13  * permission.
14  *
15  * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND ANY EXPRESS OR
16  * IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND
17  * FITNESS FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE
18  * FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING,
19  * BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS;
20  * OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT,
21  * STRICT LIABILITY, OR TOR (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
22  * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
23  *
24  **************************************************************************************************/
28 #pragma once
29 
30 #include "cutlass/fragment.h"
31 #include "cutlass/gemm/thread_multiply_add.h"
32 namespace cutlass {
33 namespace gemm {
34 
36 
38 template <typename ThreadGemmShape_,
39  typename ThreadsPerWarp_>
40 struct ThreadMultiplyAdd<ThreadGemmShape_, ThreadsPerWarp_, half, half, float> {
42  typedef Shape<1, 1, 1, 1> InstructionShape;
44  typedef ThreadGemmShape_ ThreadGemmShape;
46  typedef ThreadGemmShape AccumulatorsPerThread;
48  typedef ThreadsPerWarp_ ThreadsPerWarp;
50  typedef typename ShapeMul<ThreadGemmShape, ThreadsPerWarp>::Shape AccumulatorsPerWarp;
52  typedef half ScalarA;
54  typedef Fragment<ScalarA, AccumulatorsPerThread::kW> FragmentA;
56  typedef half ScalarB;
58  typedef Fragment<ScalarB, AccumulatorsPerThread::kH> FragmentB;
60  typedef float ScalarC;
62  typedef Fragment<ScalarC, AccumulatorsPerThread::kH * AccumulatorsPerThread::kW, 16> Accumulators;
63 
65  CUTLASS_DEVICE ThreadMultiplyAdd() {}
66 
68  CUTLASS_DEVICE void multiply_add(FragmentA const& a,
69  FragmentB const& b,
70  Accumulators const& c,
71  Accumulators& d) {
72  for (int j = 0; j < AccumulatorsPerThread::kH; ++j) {
73  for (int i = 0; i < AccumulatorsPerThread::kW; ++i) {
74  d[j * AccumulatorsPerThread::kW + i] = static_cast<ScalarC>(a[i]) * static_cast<ScalarC>(b[j]) + c[j * AccumulatorsPerThread::kW + i];
75  }
76  }
77  }
78 };
79 
81 
82 } // namespace gemm
83 } // namespace cutlass
cutlass
Definition: convert.h:33
+
cutlass::gemm::ThreadMultiplyAdd< ThreadGemmShape_, ThreadsPerWarp_, half, half, float >::ThreadMultiplyAdd
CUTLASS_DEVICE ThreadMultiplyAdd()
Ctor.
Definition: fp16_sgemm_multiply_add.h:65
+
cutlass::gemm::ThreadMultiplyAdd< ThreadGemmShape_, ThreadsPerWarp_, half, half, float >::FragmentB
Fragment< ScalarB, AccumulatorsPerThread::kH > FragmentB
The fragment for B.
Definition: fp16_sgemm_multiply_add.h:58
+
cutlass::ShapeMul::Shape
Shape< A_::kD *B_::kD, A_::kH *B_::kH, A_::kW *B_::kW, A_::kC *B_::kC > Shape
Definition: shape.h:119
+
cutlass::Fragment
A template defining Fragment Concept.
Definition: fragment.h:99
+
cutlass::gemm::ThreadMultiplyAdd< ThreadGemmShape_, ThreadsPerWarp_, half, half, float >::AccumulatorsPerWarp
ShapeMul< ThreadGemmShape, ThreadsPerWarp >::Shape AccumulatorsPerWarp
The number of accumulators per warp.
Definition: fp16_sgemm_multiply_add.h:50
+
thread_multiply_add.h
Template implementing matrix multiply-add operations on fragments.
+
cutlass::gemm::ThreadMultiplyAdd< ThreadGemmShape_, ThreadsPerWarp_, half, half, float >::ThreadGemmShape
ThreadGemmShape_ ThreadGemmShape
The shape of a thread-leveel matrix multiply accumulate.
Definition: fp16_sgemm_multiply_add.h:44
+
cutlass::gemm::ThreadMultiplyAdd< ThreadGemmShape_, ThreadsPerWarp_, half, half, float >::multiply_add
CUTLASS_DEVICE void multiply_add(FragmentA const &a, FragmentB const &b, Accumulators const &c, Accumulators &d)
Multiply : d = a*b + c.
Definition: fp16_sgemm_multiply_add.h:68
+
cutlass::gemm::ThreadMultiplyAdd< ThreadGemmShape_, ThreadsPerWarp_, half, half, float >::ScalarA
half ScalarA
The type for A. specialized to half.
Definition: fp16_sgemm_multiply_add.h:52
+
cutlass::gemm::ThreadMultiplyAdd< ThreadGemmShape_, ThreadsPerWarp_, half, half, float >::ScalarB
half ScalarB
The type for B. specialized to half.
Definition: fp16_sgemm_multiply_add.h:56
+
cutlass::gemm::ThreadMultiplyAdd< ThreadGemmShape_, ThreadsPerWarp_, half, half, float >::ThreadsPerWarp
ThreadsPerWarp_ ThreadsPerWarp
The number of threads per warp.
Definition: fp16_sgemm_multiply_add.h:48
+
cutlass::gemm::ThreadMultiplyAdd< ThreadGemmShape_, ThreadsPerWarp_, half, half, float >::FragmentA
Fragment< ScalarA, AccumulatorsPerThread::kW > FragmentA
The fragment for A.
Definition: fp16_sgemm_multiply_add.h:54
+
cutlass::gemm::ThreadMultiplyAdd< ThreadGemmShape_, ThreadsPerWarp_, half, half, float >::ScalarC
float ScalarC
The type for C and D. specialized to float.
Definition: fp16_sgemm_multiply_add.h:60
+
cutlass::Shape
A Shape implementing Layout Concept describing the dimensions of a cube.
Definition: shape.h:64
+
cutlass::gemm::ThreadMultiplyAdd< ThreadGemmShape_, ThreadsPerWarp_, half, half, float >::Accumulators
Fragment< ScalarC, AccumulatorsPerThread::kH *AccumulatorsPerThread::kW, 16 > Accumulators
The accumulators.
Definition: fp16_sgemm_multiply_add.h:62
+
cutlass::gemm::ThreadMultiplyAdd< ThreadGemmShape_, ThreadsPerWarp_, half, half, float >::AccumulatorsPerThread
ThreadGemmShape AccumulatorsPerThread
Aliased to "AccumulatorsPerThread" for compatibility. Expect to be renamed in CUTLASS v2...
Definition: fp16_sgemm_multiply_add.h:46
+
cutlass::gemm::ThreadMultiplyAdd
Template performing matrix multiply-add operation within a thread.
Definition: thread_multiply_add.h:44
+
fragment.h
Defines Fragment, a statically-sized array for storing parts of matrices within a thread&#39;s registers...
+
cutlass::gemm::ThreadMultiplyAdd< ThreadGemmShape_, ThreadsPerWarp_, half, half, float >::InstructionShape
Shape< 1, 1, 1, 1 > InstructionShape
The shape of the instruction.
Definition: fp16_sgemm_multiply_add.h:42
+
+ +
+Generated on Wed Sep 19 2018 13:58:50 for Cutlass by   +doxygen + 1.8.14 +
+ + diff --git a/docs/fp16__sgemm__traits_8h.html b/docs/fp16__sgemm__traits_8h.html new file mode 100644 index 00000000..0691fbbf --- /dev/null +++ b/docs/fp16__sgemm__traits_8h.html @@ -0,0 +1,117 @@ + + + + + + + +Cutlass: fp16_sgemm_traits.h File Reference + + + + + + + + + + +
+
+ + + + + + +
+
Cutlass +
+
CUDA Templates for Linear Algebra Subroutines and Solvers
+
+
+ + + + + + + + +
+
+ + +
+ +
+ + +
+
+
+Classes | +Namespaces
+
+
fp16_sgemm_traits.h File Reference
+
+
+ +

Defies structural properties of single-precision GEMM where any number of the input/output could be fp16 or fp32. The accumulator type stays in fp32. +More...

+
#include "cutlass/gemm/gemm.h"
+#include "cutlass/gemm/gemm_epilogue.h"
+#include "cutlass/gemm/gemm_epilogue_traits.h"
+#include "cutlass/gemm/gemm_global_tile.h"
+#include "cutlass/gemm/gemm_shared_tile.h"
+#include "cutlass/gemm/gemm_traits.h"
+#include "cutlass/gemm/fp16_sgemm_multiply_add.h"
+
+

Go to the source code of this file.

+ + + + + + +

+Classes

struct  cutlass::gemm::Fp16SgemmConfig< OutputTile_, ThreadGemmShape_, ScalarA_, ScalarB_, ScalarC_, ScalarD_, kScalarsPerLdgA_, kScalarsPerLdgB_ >
 
struct  cutlass::gemm::Fp16SgemmSgemmTraits< kLayoutA_, kLayoutB_, OutputTile_, ScalarA_, ScalarB_, ScalarC_, ScalarD_, Scalar_, EpilogueFunctor_, ThreadGemmShape_, kScalarsPerLdgA_, kScalarsPerLdgB_, Index_, GemmConfig_, GemmEpilogueTraits_ >
 
+ + + + + +

+Namespaces

 cutlass
 
 cutlass::gemm
 
+
+ +
+Generated on Wed Sep 19 2018 13:58:50 for Cutlass by   +doxygen + 1.8.14 +
+ + diff --git a/docs/fp16__sgemm__traits_8h_source.html b/docs/fp16__sgemm__traits_8h_source.html new file mode 100644 index 00000000..b5f94457 --- /dev/null +++ b/docs/fp16__sgemm__traits_8h_source.html @@ -0,0 +1,104 @@ + + + + + + + +Cutlass: fp16_sgemm_traits.h Source File + + + + + + + + + + +
+
+ + + + + + +
+
Cutlass +
+
CUDA Templates for Linear Algebra Subroutines and Solvers
+
+
+ + + + + + + + +
+
+ + +
+ +
+ + +
+
+
+
fp16_sgemm_traits.h
+
+
+Go to the documentation of this file.
1 /***************************************************************************************************
2  * Copyright (c) 2017-2018, NVIDIA CORPORATION. All rights reserved.
3  *
4  * Redistribution and use in source and binary forms, with or without modification, are permitted
5  * provided that the following conditions are met:
6  * * Redistributions of source code must retain the above copyright notice, this list of
7  * conditions and the following disclaimer.
8  * * Redistributions in binary form must reproduce the above copyright notice, this list of
9  * conditions and the following disclaimer in the documentation and/or other materials
10  * provided with the distribution.
11  * * Neither the name of the NVIDIA CORPORATION nor the names of its contributors may be used
12  * to endorse or promote products derived from this software without specific prior written
13  * permission.
14  *
15  * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND ANY EXPRESS OR
16  * IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND
17  * FITNESS FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE
18  * FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING,
19  * BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS;
20  * OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT,
21  * STRICT LIABILITY, OR TOR (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
22  * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
23  *
24  **************************************************************************************************/
29 #pragma once
30 
31 #include "cutlass/gemm/gemm.h"
32 #include "cutlass/gemm/gemm_epilogue.h"
33 #include "cutlass/gemm/gemm_epilogue_traits.h"
34 #include "cutlass/gemm/gemm_global_tile.h"
35 #include "cutlass/gemm/gemm_shared_tile.h"
36 #include "cutlass/gemm/gemm_traits.h"
37 #include "cutlass/gemm/fp16_sgemm_multiply_add.h"
38 
39 namespace cutlass {
40 namespace gemm {
41 
43 
44 template <
46  typename OutputTile_,
48  typename ThreadGemmShape_,
50  typename ScalarA_,
52  typename ScalarB_,
54  typename ScalarC_,
56  typename ScalarD_,
58  int kScalarsPerLdgA_ = 1,
60  int kScalarsPerLdgB_ = 1>
61 struct Fp16SgemmConfig : public GemmConfig<
63  ScalarA_,
65  ScalarB_,
67  ScalarC_,
69  ScalarD_,
71  OutputTile_,
73  ThreadMultiplyAdd<ThreadGemmShape_, Shape<1, 4, 8>, ScalarA_, ScalarB_, float /*for sgemm accum is float*/>,
75  kScalarsPerLdgA_,
77  kScalarsPerLdgA_,
79  4,
81  kScalarsPerLdgB_,
83  kScalarsPerLdgB_,
85  4,
87  1,
89  4,
91  1,
93  2> {};
94 
96 
97 template <
99  MatrixLayout::Kind kLayoutA_,
101  MatrixLayout::Kind kLayoutB_,
103  typename OutputTile_ = Shape<8, 128, 128>,
105  typename ScalarA_ = half,
107  typename ScalarB_ = half,
109  typename ScalarC_ = half,
111  typename ScalarD_ = half,
113  typename Scalar_ = half,
115  typename EpilogueFunctor_ = LinearScaling<Scalar_, FragmentMultiplyAdd<Scalar_, float/*accumulator type*/> >,
117  typename ThreadGemmShape_ = Shape<8, 8, 8>,
119  int kScalarsPerLdgA_ = 1,
121  int kScalarsPerLdgB_ = 1,
123  typename Index_ = int,
125  typename GemmConfig_ =
126  Fp16SgemmConfig<OutputTile_,
127  ThreadGemmShape_,
128  ScalarA_,
129  ScalarB_,
130  ScalarC_,
131  ScalarD_,
132  kScalarsPerLdgA_,
133  kScalarsPerLdgB_>,
135  typename GemmEpilogueTraits_ =
136  SimplifiedGemmEpilogueTraits<GemmConfig_, EpilogueFunctor_, Index_> >
137 struct Fp16SgemmSgemmTraits : public SimplifiedGemmTraits<
138  // The layout for A.
139  kLayoutA_,
140  // The layout for B.
141  kLayoutB_,
142  // The config.
143  GemmConfig_,
144  // The epilogue.
145  GemmEpilogue<GemmEpilogueTraits_>,
146  // The index.
147  Index_> {};
148 
150 
151 } // namespace gemm
152 } // namespace cutlass
cutlass
Definition: convert.h:33
+
gemm_global_tile.h
Defines iterators for efficiently loading and storing to global memory.
+
gemm_traits.h
Defines structural properties of complete GEMM computation.
+
gemm_epilogue.h
Implements the epilogue phase of the GEMM kernel that efficiently updates global memory with the comp...
+
gemm_shared_tile.h
Defines iterators for efficiently loading and storing tiles to and from shared memory.
+
cutlass::gemm::GemmConfig
Definition: gemm_config.h:76
+
cutlass::Shape
A Shape implementing Layout Concept describing the dimensions of a cube.
Definition: shape.h:64
+
cutlass::gemm::SimplifiedGemmEpilogueTraits
Definition: gemm_epilogue_traits.h:323
+
cutlass::gemm::Fp16SgemmConfig
Definition: fp16_sgemm_traits.h:61
+
cutlass::MatrixLayout::Kind
Kind
Enumeration defining fundamental contiguous layouts.
Definition: matrix_traits.h:159
+
fp16_sgemm_multiply_add.h
Template implementing matrix multiply-add operations on fragments.
+
cutlass::gemm::LinearScaling
Functor to compute linear combination of fragments.
Definition: linear_scaling.h:51
+
gemm.h
Implements a software-pipelined efficient GEMM.
+
gemm_epilogue_traits.h
Defines structural properties of the GEMM epilogue.
+
cutlass::gemm::Fp16SgemmSgemmTraits
Definition: fp16_sgemm_traits.h:137
+
cutlass::gemm::SimplifiedGemmTraits
Definition: gemm_traits.h:650
+
cutlass::gemm::FragmentMultiplyAdd
Definition: fragment_multiply_add.h:41
+
+ +
+Generated on Wed Sep 19 2018 13:58:50 for Cutlass by   +doxygen + 1.8.14 +
+ + diff --git a/docs/fragment_8h.html b/docs/fragment_8h.html index d97ac7b5..687dfdc8 100644 --- a/docs/fragment_8h.html +++ b/docs/fragment_8h.html @@ -83,15 +83,15 @@ $(function() {

Defines Fragment, a statically-sized array for storing parts of matrices within a thread's registers. More...

#include <assert.h>
-#include <cutlass/shape.h>
-#include <cutlass/util/cutlass_math.h>
-#include <cutlass/vector.h>
+#include "cutlass/shape.h"
+#include "cutlass/util/cutlass_math.h"
+#include "cutlass/vector.h"

Go to the source code of this file.

- + @@ -116,7 +116,7 @@ Namespaces
-Generated on Thu May 3 2018 16:36:12 for Cutlass by   +Generated on Wed Sep 19 2018 13:58:50 for Cutlass by   doxygen 1.8.14
diff --git a/docs/fragment_8h_source.html b/docs/fragment_8h_source.html index 8006bbbd..f7d23656 100644 --- a/docs/fragment_8h_source.html +++ b/docs/fragment_8h_source.html @@ -76,64 +76,66 @@ $(function() {
fragment.h
-Go to the documentation of this file.
1 /***************************************************************************************************
2  * Copyright (c) 2017-2018, NVIDIA CORPORATION. All rights reserved.
3  *
4  * Redistribution and use in source and binary forms, with or without modification, are permitted
5  * provided that the following conditions are met:
6  * * Redistributions of source code must retain the above copyright notice, this list of
7  * conditions and the following disclaimer.
8  * * Redistributions in binary form must reproduce the above copyright notice, this list of
9  * conditions and the following disclaimer in the documentation and/or other materials
10  * provided with the distribution.
11  * * Neither the name of the NVIDIA CORPORATION nor the names of its contributors may be used
12  * to endorse or promote products derived from this software without specific prior written
13  * permission.
14  *
15  * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND ANY EXPRESS OR
16  * IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND
17  * FITNESS FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE
18  * FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING,
19  * BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS;
20  * OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT,
21  * STRICT LIABILITY, OR TOR (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
22  * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
23  *
24  **************************************************************************************************/
29 #pragma once
30 
31 #include <assert.h>
32 #include <cutlass/shape.h>
33 #include <cutlass/util/cutlass_math.h>
34 #include <cutlass/vector.h>
35 
36 namespace cutlass {
37 
39 
56 
73 
75 template <int kAlignment_>
76 struct StorageType {
77  typedef uint64_t Type;
78 };
79 template <>
80 struct StorageType<4> {
81  typedef uint32_t Type;
82 };
83 template <>
84 struct StorageType<2> {
85  typedef uint16_t Type;
86 };
87 template <>
88 struct StorageType<1> {
89  typedef uint8_t Type;
90 };
91 
93 
98 template <typename Element_, int kElements_, size_t kAlignment_ = 16>
99 struct Fragment : public AlignedStruct<kAlignment_> {
101  static_assert(kAlignment_ == 16 || kAlignment_ >= sizeof(Element_), "Alignment is too small");
103  static_assert(is_pow2<kAlignment_>::value, "Alignment must be a power of two");
104 
106  typedef Fragment<Element_, kElements_> This_;
108  typedef Element_ Element;
110  static int const kElements = kElements_;
111 
113  CUTLASS_DEVICE void clear() {
114  // Avoid element-wise access for sub 32b element type
115  if (kAlignment_ >= 8 && (kElements * sizeof(Element)) % 8 == 0) {
116  uint64_t* ptr = reinterpret_cast<uint64_t*>(storage);
117  for (int i = 0; i < (kElements * sizeof(Element)) / 8; ++i) {
118  ptr[i] = uint64_t(0);
119  }
120  } else if (kAlignment_ >= 4 && (kElements * sizeof(Element)) % 4 == 0) {
121  uint32_t* ptr = reinterpret_cast<uint32_t*>(storage);
122  for (int i = 0; i < (kElements * sizeof(Element)) / 4; ++i) {
123  ptr[i] = uint32_t(0);
124  }
125  } else if (kAlignment_ >= 2 && (kElements * sizeof(Element)) % 2 == 0) {
126  uint16_t* ptr = reinterpret_cast<uint16_t*>(storage);
127  for (int i = 0; i < (kElements * sizeof(Element)) / 2; ++i) {
128  ptr[i] = uint16_t(0);
129  }
130  } else {
131  for (int i = 0; i < kElements; ++i) {
132  storage[i] = 0;
133  }
134  }
135  }
136 
138  CUTLASS_DEVICE Element& operator[](int i) {
139  assert(i < kElements_);
140  return reinterpret_cast<Element*>(storage)[i];
141  }
142 
144  CUTLASS_DEVICE Element const& operator[](int i) const {
145  assert(i < kElements_);
146  return reinterpret_cast<Element const*>(storage)[i];
147  }
148 
149  private:
151  typedef typename StorageType<kAlignment_>::Type StorageType;
152 
154  static int const kStorageCount =
155  (sizeof(Element_) * kElements_ + sizeof(StorageType) - 1) / sizeof(StorageType);
157  StorageType storage[kStorageCount];
158 
160  static_assert(sizeof(StorageType) <= kAlignment_, "StorageType is too big for given alignment");
161 };
162 
164 
169 template <typename Fragment_, typename Iterations_, typename AccessType_>
170 struct FragmentIterator {
172  typedef FragmentIterator<Fragment_, Iterations_, AccessType_> This_;
174  typedef Fragment_ Fragment;
176  typedef Iterations_ Iterations;
178  typedef AccessType_ AccessType;
179 
181  typedef typename Fragment::Element Element;
183  static int const kElementsPerAccess = (int)(sizeof(AccessType) / sizeof(Element));
185  typedef typename ShapeMul<Iterations, Shape<1, 1, 1, kElementsPerAccess> >::Shape FragmentShape;
187  typedef typename ShapeStrides<FragmentShape>::Shape Strides;
188 
190  template <typename OtherFragment_>
191  CUTLASS_DEVICE FragmentIterator(OtherFragment_& fragment, int offset = 0)
192  : pointer(reinterpret_cast<Element*>(&fragment[offset])) {
193  static_assert(OtherFragment_::kElements >= Fragment::kElements, "");
194  }
195 
197  CUTLASS_DEVICE AccessType const& at(int d, int h, int w, int c = 0) const {
198  int const imm = ComputeOffsetFromStrides<Strides>::get(d, h, w, c);
199  return reinterpret_cast<AccessType const&>(pointer[imm]);
200  }
201 
203  CUTLASS_DEVICE AccessType& at(int d, int h, int w, int c = 0) {
204  int const imm = ComputeOffsetFromStrides<Strides>::get(d, h, w, c);
205  return reinterpret_cast<AccessType&>(pointer[imm]);
206  }
207 
209  CUTLASS_DEVICE AccessType const& operator[](int i) const {
210  return reinterpret_cast<AccessType const&>(pointer[i * kElementsPerAccess]);
211  }
212 
214  CUTLASS_DEVICE AccessType& operator[](int i) {
215  return reinterpret_cast<AccessType&>(pointer[i * kElementsPerAccess]);
216  }
217 
219  CUTLASS_DEVICE bool valid(int d, int h, int w, int c) const { return true; }
220 
222  Element* pointer;
223 };
224 
226 
227 template <typename Fragment_, typename Iterations_, typename AccessType_>
228 struct FragmentConstIterator {
230  typedef FragmentIterator<Fragment_, Iterations_, AccessType_> This_;
232  typedef Fragment_ Fragment;
234  typedef Iterations_ Iterations;
236  typedef AccessType_ AccessType;
237 
239  typedef typename Fragment::Element Element;
241  static int const kElementsPerAccess = (int)(sizeof(AccessType) / sizeof(Element));
243  typedef typename ShapeMul<Iterations, Shape<1, 1, 1, kElementsPerAccess> >::Shape FragmentShape;
245  typedef typename ShapeStrides<FragmentShape>::Shape IterationsStrides;
246 
248  template <typename OtherFragment_>
249  CUTLASS_DEVICE FragmentConstIterator(OtherFragment_& fragment, int offset = 0)
250  : pointer(reinterpret_cast<Element const*>(&fragment[offset])) {
251  static_assert(OtherFragment_::kElements >= Fragment::kElements, "");
252  }
254  CUTLASS_DEVICE FragmentConstIterator(
255  FragmentIterator<Fragment_, Iterations_, AccessType_> const& rhs_)
256  : pointer(reinterpret_cast<Element const*>(rhs_.offset)) {}
257 
259  CUTLASS_DEVICE AccessType const& at(int d, int h, int w, int c = 0) const {
260  int const imm = ComputeOffsetFromStrides<IterationsStrides>::get(d, h, w, c);
261  return reinterpret_cast<AccessType const&>(pointer[imm]);
262  }
263 
265  CUTLASS_DEVICE AccessType const& operator[](int i) const {
266  return reinterpret_cast<AccessType const&>(pointer[i * kElementsPerAccess]);
267  }
268 
270  CUTLASS_DEVICE bool valid(int d, int h, int w, int c) const { return true; }
271 
273  Element const* pointer;
274 };
275 
277 
278 } // namespace cutlass
cutlass::Fragment::clear
CUTLASS_DEVICE void clear()
Clear a fragment.
Definition: fragment.h:113
+Go to the documentation of this file.
1 /***************************************************************************************************
2  * Copyright (c) 2017-2018, NVIDIA CORPORATION. All rights reserved.
3  *
4  * Redistribution and use in source and binary forms, with or without modification, are permitted
5  * provided that the following conditions are met:
6  * * Redistributions of source code must retain the above copyright notice, this list of
7  * conditions and the following disclaimer.
8  * * Redistributions in binary form must reproduce the above copyright notice, this list of
9  * conditions and the following disclaimer in the documentation and/or other materials
10  * provided with the distribution.
11  * * Neither the name of the NVIDIA CORPORATION nor the names of its contributors may be used
12  * to endorse or promote products derived from this software without specific prior written
13  * permission.
14  *
15  * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND ANY EXPRESS OR
16  * IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND
17  * FITNESS FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE
18  * FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING,
19  * BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS;
20  * OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT,
21  * STRICT LIABILITY, OR TOR (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
22  * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
23  *
24  **************************************************************************************************/
29 #pragma once
30 
31 #include <assert.h>
32 #include "cutlass/shape.h"
33 #include "cutlass/util/cutlass_math.h"
34 #include "cutlass/vector.h"
35 
36 namespace cutlass {
37 
39 
56 
73 
75 template <int alignment>
76 struct StorageType {
77  typedef uint64_t Type;
78 };
79 template <>
80 struct StorageType<4> {
81  typedef uint32_t Type;
82 };
83 template <>
84 struct StorageType<2> {
85  typedef uint16_t Type;
86 };
87 template <>
88 struct StorageType<1> {
89  typedef uint8_t Type;
90 };
91 
93 
98 template <typename Element_, int kElements_, size_t kAlignment_ = 16>
99 struct Fragment : public AlignedStruct<kAlignment_> {
101  static_assert(kAlignment_ == 16 || kAlignment_ >= sizeof(Element_), "Alignment is too small");
103  static_assert(is_pow2<kAlignment_>::value, "Alignment must be a power of two");
104 
106  typedef Fragment<Element_, kElements_> This_;
108  typedef Element_ Element;
110  static int const kElements = kElements_;
112  static int const kAlignment = kAlignment_;
113 
115  CUTLASS_HOST_DEVICE void clear() {
116  // Avoid element-wise access for sub 32b element type
117  if (kAlignment_ >= 8 && (kElements * sizeof(Element)) % 8 == 0) {
118  uint64_t* ptr = reinterpret_cast<uint64_t*>(storage);
119  for (int i = 0; i < (kElements * sizeof(Element)) / 8; ++i) {
120  ptr[i] = uint64_t(0);
121  }
122  } else if (kAlignment_ >= 4 && (kElements * sizeof(Element)) % 4 == 0) {
123  uint32_t* ptr = reinterpret_cast<uint32_t*>(storage);
124  for (int i = 0; i < (kElements * sizeof(Element)) / 4; ++i) {
125  ptr[i] = uint32_t(0);
126  }
127  } else if (kAlignment_ >= 2 && (kElements * sizeof(Element)) % 2 == 0) {
128  uint16_t* ptr = reinterpret_cast<uint16_t*>(storage);
129  for (int i = 0; i < (kElements * sizeof(Element)) / 2; ++i) {
130  ptr[i] = uint16_t(0);
131  }
132  } else {
133  for (int i = 0; i < kElements; ++i) {
134  storage[i] = 0;
135  }
136  }
137  }
138 
140  CUTLASS_HOST_DEVICE Element& operator[](int i) { return reinterpret_cast<Element*>(storage)[i]; }
141 
143  CUTLASS_HOST_DEVICE Element const& operator[](int i) const {
144  return reinterpret_cast<Element const*>(storage)[i];
145  }
146 
147  private:
149  typedef typename StorageType<kAlignment_>::Type StorageType;
150 
152  static int const kStorageCount =
153  (sizeof(Element_) * kElements_ + sizeof(StorageType) - 1) / sizeof(StorageType);
155  StorageType storage[kStorageCount];
156 
158  static_assert(sizeof(StorageType) <= kAlignment_, "StorageType is too big for given alignment");
159 };
160 
162 
167 template <typename Fragment_, typename Iterations_, typename AccessType_>
168 struct FragmentIterator {
170  typedef FragmentIterator<Fragment_, Iterations_, AccessType_> This_;
172  typedef Fragment_ Fragment;
174  typedef Iterations_ Iterations;
176  typedef AccessType_ AccessType;
177 
179  typedef typename Fragment::Element Element;
181  static int const kElementsPerAccess = (int)(sizeof(AccessType) / sizeof(Element));
183  typedef typename ShapeMul<Iterations, Shape<1, 1, 1, kElementsPerAccess> >::Shape FragmentShape;
185  typedef typename ShapeStrides<FragmentShape, kElementsPerAccess>::Shape Strides;
186 
188  template <typename OtherFragment_>
189  CUTLASS_HOST_DEVICE FragmentIterator(OtherFragment_& fragment, int offset = 0)
190  : pointer(reinterpret_cast<Element*>(&fragment[offset])) {
191  static_assert(OtherFragment_::kElements >= Fragment::kElements, "");
192  }
193 
195  CUTLASS_HOST_DEVICE AccessType const& at(int d, int h, int w, int c = 0) const {
196  int const imm = ComputeOffsetFromStrides<Strides>::get(d, h, w, c);
197  return reinterpret_cast<AccessType const&>(pointer[imm]);
198  }
199 
201  CUTLASS_HOST_DEVICE AccessType& at(int d, int h, int w, int c = 0) {
202  int const imm = ComputeOffsetFromStrides<Strides>::get(d, h, w, c);
203  return reinterpret_cast<AccessType&>(pointer[imm]);
204  }
205 
207  CUTLASS_HOST_DEVICE AccessType const& operator[](int i) const {
208  return reinterpret_cast<AccessType const&>(pointer[i * kElementsPerAccess]);
209  }
210 
212  CUTLASS_HOST_DEVICE AccessType& operator[](int i) {
213  return reinterpret_cast<AccessType&>(pointer[i * kElementsPerAccess]);
214  }
215 
217  CUTLASS_HOST_DEVICE bool valid(int d, int h, int w, int c) const { return true; }
218 
220  Element* pointer;
221 };
222 
224 
225 template <typename Fragment_, typename Iterations_, typename AccessType_>
226 struct FragmentConstIterator {
228  typedef FragmentIterator<Fragment_, Iterations_, AccessType_> This_;
230  typedef Fragment_ Fragment;
232  typedef Iterations_ Iterations;
234  typedef AccessType_ AccessType;
235 
237  typedef typename Fragment::Element Element;
239  static int const kElementsPerAccess = (int)(sizeof(AccessType) / sizeof(Element));
241  typedef typename ShapeMul<Iterations, Shape<1, 1, 1, kElementsPerAccess> >::Shape FragmentShape;
243  typedef typename ShapeStrides<FragmentShape, kElementsPerAccess>::Shape IterationsStrides;
244 
246  template <typename OtherFragment_>
247  CUTLASS_HOST_DEVICE FragmentConstIterator(OtherFragment_& fragment, int offset = 0)
248  : pointer(reinterpret_cast<Element const*>(&fragment[offset])) {
249  static_assert(OtherFragment_::kElements >= Fragment::kElements, "");
250  }
252  CUTLASS_HOST_DEVICE FragmentConstIterator(
253  FragmentIterator<Fragment_, Iterations_, AccessType_> const& rhs_)
254  : pointer(reinterpret_cast<Element const*>(rhs_.offset)) {}
255 
257  CUTLASS_HOST_DEVICE AccessType const& at(int d, int h, int w, int c = 0) const {
258  int const imm = ComputeOffsetFromStrides<IterationsStrides>::get(d, h, w, c);
259  return reinterpret_cast<AccessType const&>(pointer[imm]);
260  }
261 
263  CUTLASS_HOST_DEVICE AccessType const& operator[](int i) const {
264  return reinterpret_cast<AccessType const&>(pointer[i * kElementsPerAccess]);
265  }
266 
268  CUTLASS_HOST_DEVICE bool valid(int d, int h, int w, int c) const { return true; }
269 
271  Element const* pointer;
272 };
273 
275 
276 } // namespace cutlass
cutlass::Fragment::clear
CUTLASS_HOST_DEVICE void clear()
Clear a fragment.
Definition: fragment.h:115
+
cutlass::FragmentIterator::valid
CUTLASS_HOST_DEVICE bool valid(int d, int h, int w, int c) const
Is the iterator valid?
Definition: fragment.h:217
cutlass
Definition: convert.h:33
-
cutlass::Fragment::operator[]
CUTLASS_DEVICE Element & operator[](int i)
The accessor.
Definition: fragment.h:138
-
cutlass::FragmentIterator::at
CUTLASS_DEVICE AccessType & at(int d, int h, int w, int c=0)
The accessor.
Definition: fragment.h:203
-
cutlass::AlignedStruct
Definition: vector.h:41
-
cutlass::FragmentConstIterator
Definition: fragment.h:228
-
cutlass::FragmentConstIterator::operator[]
CUTLASS_DEVICE AccessType const & operator[](int i) const
The accessor.
Definition: fragment.h:265
-
cutlass::ShapeStrides::Shape
Shape< Shape_::kH *Shape_::kW *Shape_::kC, Shape_::kW *Shape_::kC, Shape_::kC, 1 > Shape
Definition: shape.h:155
+
cutlass::ShapeStrides::Shape
Shape< Shape_::kH *Shape_::kW *Shape_::kC, Shape_::kW *Shape_::kC, Shape_::kC, elementsPerAccess > Shape
Definition: shape.h:170
+
cutlass::AlignedStruct
Definition: vector.h:42
+
cutlass::FragmentConstIterator
Definition: fragment.h:226
+
cutlass::FragmentIterator::FragmentIterator
CUTLASS_HOST_DEVICE FragmentIterator(OtherFragment_ &fragment, int offset=0)
Ctor.
Definition: fragment.h:189
cutlass::Fragment
A template defining Fragment Concept.
Definition: fragment.h:99
-
cutlass::FragmentIterator::Element
Fragment::Element Element
The element.
Definition: fragment.h:181
-
cutlass::FragmentConstIterator::kElementsPerAccess
static int const kElementsPerAccess
The number of elements per access.
Definition: fragment.h:241
-
cutlass::FragmentIterator::Fragment
Fragment_ Fragment
The fragment.
Definition: fragment.h:174
-
cutlass::FragmentConstIterator::Fragment
Fragment_ Fragment
The fragment.
Definition: fragment.h:232
-
cutlass::FragmentIterator::operator[]
CUTLASS_DEVICE AccessType & operator[](int i)
The accessor.
Definition: fragment.h:214
-
cutlass::FragmentConstIterator::Element
Fragment::Element Element
The element.
Definition: fragment.h:239
-
cutlass::FragmentConstIterator::IterationsStrides
ShapeStrides< FragmentShape >::Shape IterationsStrides
The linear strides for iterations.
Definition: fragment.h:245
-
cutlass::FragmentConstIterator::valid
CUTLASS_DEVICE bool valid(int d, int h, int w, int c) const
Is the iterator valid?
Definition: fragment.h:270
-
cutlass::FragmentIterator::FragmentIterator
CUTLASS_DEVICE FragmentIterator(OtherFragment_ &fragment, int offset=0)
Ctor.
Definition: fragment.h:191
+
cutlass::FragmentIterator::Element
Fragment::Element Element
The element.
Definition: fragment.h:179
+
cutlass::FragmentConstIterator::kElementsPerAccess
static int const kElementsPerAccess
The number of elements per access.
Definition: fragment.h:239
+
cutlass::FragmentIterator::Fragment
Fragment_ Fragment
The fragment.
Definition: fragment.h:172
+
cutlass::FragmentConstIterator::Fragment
Fragment_ Fragment
The fragment.
Definition: fragment.h:230
+
cutlass::FragmentConstIterator::Element
Fragment::Element Element
The element.
Definition: fragment.h:237
cutlass::Fragment::This_
Fragment< Element_, kElements_ > This_
Make sure the alignment makes sense wrt the size of elements.
Definition: fragment.h:101
-
cutlass::FragmentIterator::This_
FragmentIterator< Fragment_, Iterations_, AccessType_ > This_
This class.
Definition: fragment.h:172
-
cutlass::FragmentConstIterator::FragmentShape
ShapeMul< Iterations, Shape< 1, 1, 1, kElementsPerAccess > >::Shape FragmentShape
The shape of the the fragment.
Definition: fragment.h:243
+
cutlass::FragmentIterator::This_
FragmentIterator< Fragment_, Iterations_, AccessType_ > This_
This class.
Definition: fragment.h:170
+
cutlass::FragmentConstIterator::FragmentShape
ShapeMul< Iterations, Shape< 1, 1, 1, kElementsPerAccess > >::Shape FragmentShape
The shape of the the fragment.
Definition: fragment.h:241
cutlass_math.h
Math utilities.
cutlass::StorageType
Definition: fragment.h:76
cutlass::StorageType< 4 >::Type
uint32_t Type
Definition: fragment.h:81
cutlass::StorageType< 1 >::Type
uint8_t Type
Definition: fragment.h:89
-
cutlass::ComputeOffsetFromStrides::get
static CUTLASS_DEVICE int get(int d, int h, int w, int c)
Definition: shape.h:211
-
cutlass::FragmentIterator::pointer
Element * pointer
The pointer.
Definition: fragment.h:222
-
cutlass::FragmentConstIterator::AccessType
AccessType_ AccessType
The access type.
Definition: fragment.h:236
+
cutlass::FragmentIterator::pointer
Element * pointer
The pointer.
Definition: fragment.h:220
+
cutlass::Fragment::operator[]
CUTLASS_HOST_DEVICE Element const & operator[](int i) const
The accessor.
Definition: fragment.h:143
+
cutlass::FragmentConstIterator::AccessType
AccessType_ AccessType
The access type.
Definition: fragment.h:234
+
cutlass::FragmentConstIterator::IterationsStrides
ShapeStrides< FragmentShape, kElementsPerAccess >::Shape IterationsStrides
The linear strides for iterations.
Definition: fragment.h:243
cutlass::ShapeMul
Definition: shape.h:118
-
cutlass::FragmentIterator::FragmentShape
ShapeMul< Iterations, Shape< 1, 1, 1, kElementsPerAccess > >::Shape FragmentShape
The shape of the the fragment.
Definition: fragment.h:185
-
cutlass::FragmentIterator
A template defining Fragment Iterator Concept.
Definition: fragment.h:170
+
cutlass::FragmentIterator::FragmentShape
ShapeMul< Iterations, Shape< 1, 1, 1, kElementsPerAccess > >::Shape FragmentShape
The shape of the the fragment.
Definition: fragment.h:183
+
cutlass::FragmentConstIterator::FragmentConstIterator
CUTLASS_HOST_DEVICE FragmentConstIterator(OtherFragment_ &fragment, int offset=0)
Ctor.
Definition: fragment.h:247
+
cutlass::FragmentIterator
A template defining Fragment Iterator Concept.
Definition: fragment.h:168
cutlass::Fragment::kElements
static int const kElements
The number of elements.
Definition: fragment.h:110
-
cutlass::Fragment::operator[]
CUTLASS_DEVICE Element const & operator[](int i) const
The accessor.
Definition: fragment.h:144
-
cutlass::FragmentConstIterator::Iterations
Iterations_ Iterations
The number of iterations.
Definition: fragment.h:234
-
static_assert
#define static_assert(__e, __m)
Definition: platform.h:145
-
cutlass::FragmentIterator::Iterations
Iterations_ Iterations
The number of iterations.
Definition: fragment.h:176
+
CUTLASS_HOST_DEVICE
#define CUTLASS_HOST_DEVICE
Definition: cutlass.h:46
+
cutlass::FragmentConstIterator::Iterations
Iterations_ Iterations
The number of iterations.
Definition: fragment.h:232
+
cutlass::FragmentIterator::at
CUTLASS_HOST_DEVICE AccessType const & at(int d, int h, int w, int c=0) const
The accessor.
Definition: fragment.h:195
+
static_assert
#define static_assert(__e, __m)
Definition: platform.h:153
+
cutlass::FragmentIterator::Iterations
Iterations_ Iterations
The number of iterations.
Definition: fragment.h:174
+
cutlass::FragmentConstIterator::FragmentConstIterator
CUTLASS_HOST_DEVICE FragmentConstIterator(FragmentIterator< Fragment_, Iterations_, AccessType_ > const &rhs_)
Create from non-constant FragmentIterator.
Definition: fragment.h:252
cutlass::Shape
A Shape implementing Layout Concept describing the dimensions of a cube.
Definition: shape.h:64
-
cutlass::FragmentConstIterator::at
CUTLASS_DEVICE AccessType const & at(int d, int h, int w, int c=0) const
The accessor.
Definition: fragment.h:259
cutlass::Fragment::Element
Element_ Element
The element.
Definition: fragment.h:108
-
cutlass::FragmentConstIterator::This_
FragmentIterator< Fragment_, Iterations_, AccessType_ > This_
This class.
Definition: fragment.h:230
-
cutlass::FragmentIterator::operator[]
CUTLASS_DEVICE AccessType const & operator[](int i) const
The accessor.
Definition: fragment.h:209
+
cutlass::FragmentConstIterator::This_
FragmentIterator< Fragment_, Iterations_, AccessType_ > This_
This class.
Definition: fragment.h:228
+
cutlass::FragmentConstIterator::operator[]
CUTLASS_HOST_DEVICE AccessType const & operator[](int i) const
The accessor.
Definition: fragment.h:263
+
cutlass::Fragment::operator[]
CUTLASS_HOST_DEVICE Element & operator[](int i)
The accessor.
Definition: fragment.h:140
+
cutlass::FragmentIterator::operator[]
CUTLASS_HOST_DEVICE AccessType const & operator[](int i) const
The accessor.
Definition: fragment.h:207
cutlass::StorageType< 2 >::Type
uint16_t Type
Definition: fragment.h:85
vector.h
Defines a 1D vector of elements held in the registers of each thread.
-
cutlass::FragmentConstIterator::FragmentConstIterator
CUTLASS_DEVICE FragmentConstIterator(FragmentIterator< Fragment_, Iterations_, AccessType_ > const &rhs_)
Create from non-constant FragmentIterator.
Definition: fragment.h:254
-
cutlass::FragmentIterator::kElementsPerAccess
static int const kElementsPerAccess
The number of elements per access.
Definition: fragment.h:183
-
cutlass::FragmentIterator::Strides
ShapeStrides< FragmentShape >::Shape Strides
The linear strides for iterations.
Definition: fragment.h:187
+
cutlass::StorageType::Type
uint64_t Type
Definition: fragment.h:77
+
cutlass::FragmentConstIterator::valid
CUTLASS_HOST_DEVICE bool valid(int d, int h, int w, int c) const
Is the iterator valid?
Definition: fragment.h:268
+
cutlass::FragmentIterator::Strides
ShapeStrides< FragmentShape, kElementsPerAccess >::Shape Strides
The linear strides for iterations.
Definition: fragment.h:185
+
cutlass::ComputeOffsetFromStrides::get
static CUTLASS_HOST_DEVICE int get(int d, int h, int w, int c)
Definition: shape.h:199
+
cutlass::FragmentIterator::operator[]
CUTLASS_HOST_DEVICE AccessType & operator[](int i)
The accessor.
Definition: fragment.h:212
+
cutlass::FragmentIterator::at
CUTLASS_HOST_DEVICE AccessType & at(int d, int h, int w, int c=0)
The accessor.
Definition: fragment.h:201
+
cutlass::FragmentIterator::kElementsPerAccess
static int const kElementsPerAccess
The number of elements per access.
Definition: fragment.h:181
shape.h
Defines Shape implementing the Layout concept for representing a 4D hypercube of objects.
-
cutlass::FragmentIterator::AccessType
AccessType_ AccessType
The access type.
Definition: fragment.h:178
-
cutlass::FragmentIterator::valid
CUTLASS_DEVICE bool valid(int d, int h, int w, int c) const
Is the iterator valid?
Definition: fragment.h:219
-
cutlass::StorageType::Type
uint64_t Type
Definition: fragment.h:77
+
cutlass::FragmentIterator::AccessType
AccessType_ AccessType
The access type.
Definition: fragment.h:176
+
cutlass::Fragment::kAlignment
static int const kAlignment
Alignment.
Definition: fragment.h:112
cutlass::is_pow2
Definition: cutlass_math.h:45
-
cutlass::FragmentConstIterator::FragmentConstIterator
CUTLASS_DEVICE FragmentConstIterator(OtherFragment_ &fragment, int offset=0)
Ctor.
Definition: fragment.h:249
-
cutlass::FragmentIterator::at
CUTLASS_DEVICE AccessType const & at(int d, int h, int w, int c=0) const
The accessor.
Definition: fragment.h:197
-
cutlass::FragmentConstIterator::pointer
Element const * pointer
The pointer.
Definition: fragment.h:273
+
cutlass::FragmentConstIterator::at
CUTLASS_HOST_DEVICE AccessType const & at(int d, int h, int w, int c=0) const
The accessor.
Definition: fragment.h:257
+
cutlass::FragmentConstIterator::pointer
Element const * pointer
The pointer.
Definition: fragment.h:271
-Generated on Thu May 3 2018 16:36:12 for Cutlass by   +Generated on Wed Sep 19 2018 13:58:50 for Cutlass by   doxygen 1.8.14
diff --git a/docs/fragment__multiply__add_8h.html b/docs/fragment__multiply__add_8h.html index 59a94dfd..107cfee7 100644 --- a/docs/fragment__multiply__add_8h.html +++ b/docs/fragment__multiply__add_8h.html @@ -82,15 +82,15 @@ $(function() {

Defines multiply-add operations on fragments within a thread. More...

-
#include <cutlass/fragment.h>
+
#include "cutlass/fragment.h"

Go to the source code of this file.

Classes

struct  cutlass::StorageType< kAlignment_ >
struct  cutlass::StorageType< alignment >
 
struct  cutlass::StorageType< 4 >
 
- + - +

Classes

struct  cutlass::gemm::FragmentMultiplyAdd< Scalar_ >
struct  cutlass::gemm::FragmentMultiplyAdd< ScalarAlphaBeta_, ScalarAccum_, fragMul2 >
 
struct  cutlass::gemm::FragmentMultiplyAdd< half >
struct  cutlass::gemm::FragmentMultiplyAdd< half, half, true >
 
-Generated on Thu May 3 2018 16:36:12 for Cutlass by   +Generated on Wed Sep 19 2018 13:58:50 for Cutlass by   doxygen 1.8.14
diff --git a/docs/fragment__multiply__add_8h_source.html b/docs/fragment__multiply__add_8h_source.html index 9b453fd9..1d4c4f7f 100644 --- a/docs/fragment__multiply__add_8h_source.html +++ b/docs/fragment__multiply__add_8h_source.html @@ -76,28 +76,26 @@ $(function() {
fragment_multiply_add.h
-Go to the documentation of this file.
1 /***************************************************************************************************
2  * Copyright (c) 2017-2018, NVIDIA CORPORATION. All rights reserved.
3  *
4  * Redistribution and use in source and binary forms, with or without modification, are permitted
5  * provided that the following conditions are met:
6  * * Redistributions of source code must retain the above copyright notice, this list of
7  * conditions and the following disclaimer.
8  * * Redistributions in binary form must reproduce the above copyright notice, this list of
9  * conditions and the following disclaimer in the documentation and/or other materials
10  * provided with the distribution.
11  * * Neither the name of the NVIDIA CORPORATION nor the names of its contributors may be used
12  * to endorse or promote products derived from this software without specific prior written
13  * permission.
14  *
15  * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND ANY EXPRESS OR
16  * IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND
17  * FITNESS FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE
18  * FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING,
19  * BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS;
20  * OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT,
21  * STRICT LIABILITY, OR TOR (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
22  * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
23  *
24  **************************************************************************************************/
28 #pragma once
29 
30 #include <cutlass/fragment.h>
31 
32 namespace cutlass {
33 namespace gemm {
34 
36 
37 template <typename Scalar_>
38 struct FragmentMultiplyAdd {
40  typedef Shape<1, 1, 1, 1> InstructionShape;
42  typedef Scalar_ ScalarA;
44  typedef Scalar_ ScalarB;
46  typedef Scalar_ ScalarC;
47 
49  CUTLASS_DEVICE FragmentMultiplyAdd() {}
50 
52  template <typename Fragment_>
53  CUTLASS_DEVICE void multiply(Scalar_ a, Fragment_ const& b, Fragment_& d) {
54  for (int j = 0; j < Fragment_::kElements; ++j) {
55  d[j] = a * b[j];
56  }
57  }
58 
60  template <typename Fragment_>
61  CUTLASS_DEVICE void multiply_add(Scalar_ a,
62  Fragment_ const& b,
63  Fragment_ const& c,
64  Fragment_& d) {
65  for (int j = 0; j < Fragment_::kElements; ++j) {
66  d[j] = a * b[j] + c[j];
67  }
68  }
69 };
70 
72 
73 #if !defined(__CUDACC_RTC__) || defined(CUTLASS_NVRTC_HAS_FP16)
74 template <>
75 struct FragmentMultiplyAdd<half> {
77  typedef Shape<1, 1, 1, 1> InstructionShape;
79  typedef half ScalarA;
81  typedef half ScalarB;
83  typedef half ScalarC;
84 
86  CUTLASS_DEVICE FragmentMultiplyAdd() {}
87 
89  template <typename Fragment_>
90  CUTLASS_DEVICE void multiply(half a, Fragment_ const& b, Fragment_& d) {
91 #if defined(__CUDACC__) && __CUDA_ARCH__ >= 530
92  // The input.
93  __half2 const* b_half2 = reinterpret_cast<__half2 const*>(&b[0]);
94  // The output.
95  __half2* d_half2 = reinterpret_cast<__half2*>(&d[0]);
96 
97  // Assemble a half2 from a.
98  __half2 const a_half2 = __half2half2(a);
99 
100  for (int i = 0; i < Fragment_::kElements / 2; ++i) {
101  d_half2[i] = __hmul2(a_half2, b_half2[i]);
102  }
103 #endif
104  }
105 
107  template <typename Fragment_>
108  CUTLASS_DEVICE void multiply_add(half a, Fragment_ const& b, Fragment_ const& c, Fragment_& d) {
109 #if defined(__CUDACC__) && __CUDA_ARCH__ >= 530
110  // The inputs.
111  __half2 const* b_half2 = reinterpret_cast<__half2 const*>(&b[0]);
112  __half2 const* c_half2 = reinterpret_cast<__half2 const*>(&c[0]);
113  // The output.
114  __half2* d_half2 = reinterpret_cast<__half2*>(&d[0]);
115 
116  // Assemble a half2 from a.
117  __half2 const a_half2 = __half2half2(a);
118 
119  for (int i = 0; i < Fragment_::kElements / 2; ++i) {
120  d_half2[i] = __hfma2(a_half2, b_half2[i], c_half2[i]);
121  }
122 #endif
123  }
124 };
125 
126 #endif
127 
129 
130 } // namespace gemm
131 } // namespace cutlass
cutlass::gemm::FragmentMultiplyAdd::ScalarB
Scalar_ ScalarB
The type for B.
Definition: fragment_multiply_add.h:44
+Go to the documentation of this file.
1 /***************************************************************************************************
2  * Copyright (c) 2017-2018, NVIDIA CORPORATION. All rights reserved.
3  *
4  * Redistribution and use in source and binary forms, with or without modification, are permitted
5  * provided that the following conditions are met:
6  * * Redistributions of source code must retain the above copyright notice, this list of
7  * conditions and the following disclaimer.
8  * * Redistributions in binary form must reproduce the above copyright notice, this list of
9  * conditions and the following disclaimer in the documentation and/or other materials
10  * provided with the distribution.
11  * * Neither the name of the NVIDIA CORPORATION nor the names of its contributors may be used
12  * to endorse or promote products derived from this software without specific prior written
13  * permission.
14  *
15  * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND ANY EXPRESS OR
16  * IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND
17  * FITNESS FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE
18  * FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING,
19  * BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS;
20  * OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT,
21  * STRICT LIABILITY, OR TOR (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
22  * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
23  *
24  **************************************************************************************************/
28 #pragma once
29 
30 #include "cutlass/fragment.h"
31 
32 namespace cutlass {
33 namespace gemm {
34 
36 
37 template < typename ScalarAlphaBeta_,
38  typename ScalarAccum_,
39  bool fragMul2 = true /*number of element per fragment is multiple of 2*/
40 >
41 struct FragmentMultiplyAdd {
43  typedef Shape<1, 1, 1, 1> InstructionShape;
45  typedef ScalarAlphaBeta_ ScalarAlphaBeta;
47  typedef ScalarAccum_ ScalarAccum;
48 
50  CUTLASS_DEVICE FragmentMultiplyAdd() {}
51 
53  template <typename FragmentB_, typename FragmentCd_>
54  CUTLASS_DEVICE void multiply(ScalarAlphaBeta a, FragmentB_ const& b, FragmentCd_& d) {
55 #if defined(__CUDACC__) && __CUDA_ARCH__ >= 530
56  int const kReduction = FragmentB_::kElements / FragmentCd_::kElements;
57  for (int j = 0; j < FragmentCd_::kElements; ++j) {
58  d[j] = b[j * kReduction + 0];
59  for (int k = 1; k < kReduction; ++k) {
60  d[j] += b[j * kReduction + k];
61  }
62  d[j] = a * ScalarAlphaBeta(d[j]);
63  }
64 #endif
65  }
66 
68  template <typename FragmentB_, typename FragmentCd_>
69  CUTLASS_DEVICE void multiply_add(ScalarAlphaBeta a,
70  FragmentB_ const& b,
71  FragmentCd_ const& c,
72  FragmentCd_& d) {
73 #if defined(__CUDACC__) && __CUDA_ARCH__ >= 530
74  int const kReduction = FragmentB_::kElements / FragmentCd_::kElements;
75  for (int j = 0; j < FragmentCd_::kElements; ++j) {
76  d[j] = b[j * kReduction + 0];
77  for (int k = 1; k < kReduction; ++k) {
78  d[j] += b[j * kReduction + k];
79  }
80  d[j] = a * ScalarAlphaBeta(d[j]) + ScalarAlphaBeta(c[j]);
81  }
82 #endif
83  }
84 };
85 
87 
88 #if !defined(__CUDACC_RTC__) || defined(CUTLASS_NVRTC_HAS_FP16)
89 template <>
90 struct FragmentMultiplyAdd<half, half, true> {
92  typedef Shape<1, 1, 1, 1> InstructionShape;
94  typedef half ScalarAlphaBeta;
96  typedef half ScalarAccum;
97 
99  CUTLASS_DEVICE FragmentMultiplyAdd() {}
100 
102  template <typename FragmentB_, typename FragmentCd_>
103  CUTLASS_DEVICE void multiply(half a, FragmentB_ const& b, FragmentCd_& d) {
104 #if defined(__CUDACC__) && __CUDA_ARCH__ >= 530
105  // The input.
106  __half2 const* b_half2 = reinterpret_cast<__half2 const*>(&b[0]);
107  // The output.
108  __half2* d_half2 = reinterpret_cast<__half2*>(&d[0]);
109 
110  // Assemble a half2 from a.
111  __half2 const a_half2 = __half2half2(a);
112 
113  int const kReduction = (FragmentB_::kElements / FragmentCd_::kElements);
114 
115  for (int j = 0; j < FragmentCd_::kElements / 2; ++j) {
116  d_half2[j] = __hmul2(a_half2, b_half2[j * kReduction + 0]);
117 
118  for (int k = 1; k < kReduction; ++k) {
119  d_half2[j] = __hfma2(a_half2, b_half2[j * kReduction + k], d_half2[j]);
120  }
121  }
122 #endif
123  }
124 
125 
127  template <typename FragmentB_, typename FragmentCd_>
128  CUTLASS_DEVICE void multiply_add(half a,
129  FragmentB_ const& b,
130  FragmentCd_ const& c,
131  FragmentCd_& d) {
132 #if defined(__CUDACC__) && __CUDA_ARCH__ >= 530
133  // The inputs.
134  __half2 const* b_half2 = reinterpret_cast<__half2 const*>(&b[0]);
135  __half2 const* c_half2 = reinterpret_cast<__half2 const*>(&c[0]);
136  // The output.
137  __half2* d_half2 = reinterpret_cast<__half2*>(&d[0]);
138 
139  // Assemble a half2 from a.
140  __half2 const a_half2 = __half2half2(a);
141 
142  int const kReduction = (FragmentB_::kElements / FragmentCd_::kElements);
143  for (int j = 0; j < FragmentCd_::kElements / 2; ++j) {
144  d_half2[j] = __hfma2(a_half2, b_half2[j * kReduction + 0], c_half2[j]);
145 
146  for (int k = 1; k < kReduction; ++k) {
147  d_half2[j] = __hfma2(a_half2, b_half2[j * kReduction + k], d_half2[j]);
148  }
149  }
150 #endif
151  }
152 };
153 
154 #endif
155 
157 
158 } // namespace gemm
159 } // namespace cutlass
cutlass::gemm::FragmentMultiplyAdd::multiply
CUTLASS_DEVICE void multiply(ScalarAlphaBeta a, FragmentB_ const &b, FragmentCd_ &d)
Multiply : d = a*b.
Definition: fragment_multiply_add.h:54
+
cutlass::gemm::FragmentMultiplyAdd< half, half, true >::InstructionShape
Shape< 1, 1, 1, 1 > InstructionShape
The shape of the instruction.
Definition: fragment_multiply_add.h:92
cutlass
Definition: convert.h:33
-
cutlass::gemm::FragmentMultiplyAdd::multiply
CUTLASS_DEVICE void multiply(Scalar_ a, Fragment_ const &b, Fragment_ &d)
Multiply : d = a*b.
Definition: fragment_multiply_add.h:53
-
cutlass::gemm::FragmentMultiplyAdd< half >::ScalarA
half ScalarA
The type for A.
Definition: fragment_multiply_add.h:79
-
cutlass::gemm::FragmentMultiplyAdd< half >::FragmentMultiplyAdd
CUTLASS_DEVICE FragmentMultiplyAdd()
Ctor.
Definition: fragment_multiply_add.h:86
-
cutlass::gemm::FragmentMultiplyAdd::multiply_add
CUTLASS_DEVICE void multiply_add(Scalar_ a, Fragment_ const &b, Fragment_ const &c, Fragment_ &d)
Multiply : d = a*b + c.
Definition: fragment_multiply_add.h:61
-
cutlass::gemm::FragmentMultiplyAdd< half >::ScalarC
half ScalarC
The type for C and D.
Definition: fragment_multiply_add.h:83
-
cutlass::gemm::FragmentMultiplyAdd< half >::multiply_add
CUTLASS_DEVICE void multiply_add(half a, Fragment_ const &b, Fragment_ const &c, Fragment_ &d)
Multiply : d = a*b + c.
Definition: fragment_multiply_add.h:108
+
cutlass::gemm::FragmentMultiplyAdd< half, half, true >::ScalarAlphaBeta
half ScalarAlphaBeta
The type for alpha and beta.
Definition: fragment_multiply_add.h:94
+
cutlass::gemm::FragmentMultiplyAdd::FragmentMultiplyAdd
CUTLASS_DEVICE FragmentMultiplyAdd()
Ctor.
Definition: fragment_multiply_add.h:50
+
cutlass::gemm::FragmentMultiplyAdd< half, half, true >::FragmentMultiplyAdd
CUTLASS_DEVICE FragmentMultiplyAdd()
Ctor.
Definition: fragment_multiply_add.h:99
+
cutlass::gemm::FragmentMultiplyAdd< half, half, true >::multiply
CUTLASS_DEVICE void multiply(half a, FragmentB_ const &b, FragmentCd_ &d)
Multiply : d = a*b.
Definition: fragment_multiply_add.h:103
+
cutlass::gemm::FragmentMultiplyAdd::ScalarAccum
ScalarAccum_ ScalarAccum
The type for accumlator.
Definition: fragment_multiply_add.h:47
cutlass::Shape
A Shape implementing Layout Concept describing the dimensions of a cube.
Definition: shape.h:64
-
cutlass::gemm::FragmentMultiplyAdd::InstructionShape
Shape< 1, 1, 1, 1 > InstructionShape
The shape of the instruction.
Definition: fragment_multiply_add.h:40
-
cutlass::gemm::FragmentMultiplyAdd::ScalarC
Scalar_ ScalarC
The type for C and D.
Definition: fragment_multiply_add.h:46
-
cutlass::gemm::FragmentMultiplyAdd::ScalarA
Scalar_ ScalarA
The type for A.
Definition: fragment_multiply_add.h:42
-
cutlass::gemm::FragmentMultiplyAdd::FragmentMultiplyAdd
CUTLASS_DEVICE FragmentMultiplyAdd()
Ctor.
Definition: fragment_multiply_add.h:49
+
cutlass::gemm::FragmentMultiplyAdd::ScalarAlphaBeta
ScalarAlphaBeta_ ScalarAlphaBeta
The type for alpha and beta.
Definition: fragment_multiply_add.h:45
+
cutlass::gemm::FragmentMultiplyAdd< half, half, true >::multiply_add
CUTLASS_DEVICE void multiply_add(half a, FragmentB_ const &b, FragmentCd_ const &c, FragmentCd_ &d)
Multiply : d = a*b + c.
Definition: fragment_multiply_add.h:128
+
cutlass::gemm::FragmentMultiplyAdd::InstructionShape
Shape< 1, 1, 1, 1 > InstructionShape
The shape of the instruction.
Definition: fragment_multiply_add.h:43
fragment.h
Defines Fragment, a statically-sized array for storing parts of matrices within a thread&#39;s registers...
-
cutlass::gemm::FragmentMultiplyAdd< half >::multiply
CUTLASS_DEVICE void multiply(half a, Fragment_ const &b, Fragment_ &d)
Multiply : d = a*b.
Definition: fragment_multiply_add.h:90
-
cutlass::gemm::FragmentMultiplyAdd< half >::InstructionShape
Shape< 1, 1, 1, 1 > InstructionShape
The shape of the instruction.
Definition: fragment_multiply_add.h:77
-
cutlass::gemm::FragmentMultiplyAdd< half >::ScalarB
half ScalarB
The type for B.
Definition: fragment_multiply_add.h:81
-
cutlass::gemm::FragmentMultiplyAdd
Definition: fragment_multiply_add.h:38
+
cutlass::gemm::FragmentMultiplyAdd< half, half, true >::ScalarAccum
half ScalarAccum
The type for accumlator.
Definition: fragment_multiply_add.h:96
+
cutlass::gemm::FragmentMultiplyAdd::multiply_add
CUTLASS_DEVICE void multiply_add(ScalarAlphaBeta a, FragmentB_ const &b, FragmentCd_ const &c, FragmentCd_ &d)
Multiply : d = a*b + c.
Definition: fragment_multiply_add.h:69
+
cutlass::gemm::FragmentMultiplyAdd
Definition: fragment_multiply_add.h:41
-Generated on Thu May 3 2018 16:36:12 for Cutlass by   +Generated on Wed Sep 19 2018 13:58:50 for Cutlass by   doxygen 1.8.14
diff --git a/docs/functions.html b/docs/functions.html index e6b156fb..bdde612a 100644 --- a/docs/functions.html +++ b/docs/functions.html @@ -71,77 +71,101 @@ $(function() {
Here is a list of all class members with links to the classes they belong to:

- a -

-Generated on Thu May 3 2018 16:36:15 for Cutlass by   +Generated on Wed Sep 19 2018 13:58:53 for Cutlass by   doxygen 1.8.14
diff --git a/docs/functions_0x7e.html b/docs/functions_0x7e.html index 41aa664c..0cb0e345 100644 --- a/docs/functions_0x7e.html +++ b/docs/functions_0x7e.html @@ -78,7 +78,7 @@ $(function() {
-Generated on Thu May 3 2018 16:36:15 for Cutlass by   +Generated on Wed Sep 19 2018 13:58:53 for Cutlass by   doxygen 1.8.14
diff --git a/docs/functions_b.html b/docs/functions_b.html index 79038aa1..35c2018b 100644 --- a/docs/functions_b.html +++ b/docs/functions_b.html @@ -71,42 +71,82 @@ $(function() {
Here is a list of all class members with links to the classes they belong to:

- b -

-Generated on Thu May 3 2018 16:36:15 for Cutlass by   +Generated on Wed Sep 19 2018 13:58:53 for Cutlass by   doxygen 1.8.14
diff --git a/docs/functions_enum.html b/docs/functions_enum.html index b710de0f..df8ae39f 100644 --- a/docs/functions_enum.html +++ b/docs/functions_enum.html @@ -70,18 +70,20 @@ $(function() {
 
-Generated on Thu May 3 2018 16:36:15 for Cutlass by   +Generated on Wed Sep 19 2018 13:58:53 for Cutlass by   doxygen 1.8.14
diff --git a/docs/functions_eval.html b/docs/functions_eval.html index 40c01ec8..b5312933 100644 --- a/docs/functions_eval.html +++ b/docs/functions_eval.html @@ -77,6 +77,13 @@ $(function() { +

- b -

+ +

- k -

@@ -139,6 +160,13 @@ $(function() { +

- o -

+ +

- v -

-Generated on Thu May 3 2018 16:36:15 for Cutlass by   +Generated on Wed Sep 19 2018 13:58:53 for Cutlass by   doxygen 1.8.14
diff --git a/docs/functions_func_g.html b/docs/functions_func_g.html index b3023736..b6258df4 100644 --- a/docs/functions_func_g.html +++ b/docs/functions_func_g.html @@ -74,47 +74,73 @@ $(function() {
  • Gemm() : cutlass::gemm::Gemm< GemmTraits_ >
    • +
  • GemmCoord() +: cutlass::gemm::GemmCoord +
    • +
  • GemmDesc() +: cutlass::gemm::GemmDesc< AType_, BType_, CType_, DType_, SType_, Index_ > +
  • GemmEpilogue() -: cutlass::gemm::GemmEpilogue< GemmEpilogueTraits_ > +: cutlass::gemm::GemmEpilogue< GemmEpilogueTraits_ >
  • GemmGlobalIteratorAb() -: cutlass::gemm::GemmGlobalIteratorAb< TileTraits_, Index_ > +: cutlass::gemm::GemmGlobalIteratorAb< TileTraits_, Index_ >
  • GemmGlobalIteratorCd() -: cutlass::gemm::GemmGlobalIteratorCd< TileTraits_, Index_ > +: cutlass::gemm::GemmGlobalIteratorCd< TileTraits_, Index_ >
  • get() -: cutlass::ComputeOffsetFromShape< Shape_ > -, cutlass::ComputeOffsetFromShape< Shape< 1, kSh_, kSw_, 1 > > -, cutlass::ComputeOffsetFromShape< Shape< 1, kSh_, kSw_, kSc_ > > -, cutlass::ComputeOffsetFromStrides< Strides_ > -, cutlass::ComputeOffsetFromStrides< Shape< 1, S_h_, S_w_, 1 > > -, cutlass::ComputeOffsetFromStrides< Shape< 1, S_h_, S_w_, S_c_ > > +: cutlass::ComputeOffsetFromShape< Shape_ > +, cutlass::ComputeOffsetFromStrides< Strides_ > , cutlass::ComputeThreadOffsetFromStrides< Threads_, Strides_ > , cutlass::ComputeThreadOffsetFromStrides< Shape< 1, T_h_, T_w_, 1 >, Shape< 1, S_h_, S_w_, 1 > > , cutlass::ComputeThreadOffsetFromStrides< Shape< 1, T_h_, T_w_, T_c_ >, Shape< 1, S_h_, S_w_, S_c_ > > +, cutlass::detail::ScalarOrPointer< Scalar_ > , cutlass::platform::unique_ptr< T, Deleter > , cutlass::PredicateVector< kPredicates_, kPredicatesPerByte_, kPredicateStart_ >::Iterator
    • +
  • get_batch_id() +: cutlass::gemm::ColumnMajorBlockSwizzle< groupCols, swDirection > +, cutlass::gemm::IdentityBlockSwizzle +, cutlass::gemm::RowMajorBlockSwizzle< groupRows, swDirection > +
  • get_deleter() -: cutlass::platform::unique_ptr< T, Deleter > +: cutlass::platform::unique_ptr< T, Deleter > +
    • +
  • get_grid_layout() +: cutlass::gemm::ColumnMajorBlockSwizzle< groupCols, swDirection > +, cutlass::gemm::IdentityBlockSwizzle +, cutlass::gemm::RowMajorBlockSwizzle< groupRows, swDirection > +
    • +
  • get_pointer_offset() +: cutlass::TensorRefBatchStrided< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ > +
    • +
  • get_ptr() +: cutlass::detail::ScalarOrPointer< Scalar_ > +
    • +
  • get_scalar() +: cutlass::detail::ScalarOrPointer< Scalar_ > +
    • +
  • get_threadblock_offset() +: cutlass::gemm::ColumnMajorBlockSwizzle< groupCols, swDirection > +, cutlass::gemm::IdentityBlockSwizzle +, cutlass::gemm::RowMajorBlockSwizzle< groupRows, swDirection >
  • GlobalLoadStream() -: cutlass::gemm::GemmTraits< GemmConfig_, GlobalLoadStreamA_, GlobalLoadStreamB_, SharedLoadStreamA_, SharedLoadStreamB_, Epilogue_, BlockSwizzle_, Index_, ClearAccumulators_ >::GlobalLoadStream -, cutlass::gemm::GlobalLoadStream< LoadIterator_, StoreIterator_, Transformer_ > +: cutlass::gemm::GlobalLoadStream< Operand, LoadIterator_, StoreIterator_, Transformer_ >
    • -
  • GlobalLoadStreamBase() -: cutlass::gemm::GlobalLoadStreamBase< LoadIterator_, StoreIterator_, Transformer_ > +
  • GlobalLoadStreamPair() +: cutlass::gemm::GlobalLoadStreamPair< StreamA_, StreamB_, kResidueInProlog_ >
  • good() -: cutlass::TensorRef< Storage_, Rank_ > -, cutlass::TensorView< T > +: cutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ > +, cutlass::TensorRef< Storage_, Rank_, MapFunc_, 1, Index_, LongIndex_ >
    • -Generated on Thu May 3 2018 16:36:15 for Cutlass by   +Generated on Wed Sep 19 2018 13:58:53 for Cutlass by   doxygen 1.8.14
    diff --git a/docs/functions_func_h.html b/docs/functions_func_h.html index 7eb85aac..184c449e 100644 --- a/docs/functions_func_h.html +++ b/docs/functions_func_h.html @@ -71,14 +71,23 @@ $(function() {  

    - h -

    -Generated on Thu May 3 2018 16:36:15 for Cutlass by   +Generated on Wed Sep 19 2018 13:58:53 for Cutlass by   doxygen 1.8.14
    diff --git a/docs/functions_func_i.html b/docs/functions_func_i.html index 16cfdc51..7c8680ab 100644 --- a/docs/functions_func_i.html +++ b/docs/functions_func_i.html @@ -72,90 +72,103 @@ $(function() {

    - i -

    -Generated on Thu May 3 2018 16:36:15 for Cutlass by   +Generated on Wed Sep 19 2018 13:58:53 for Cutlass by   doxygen 1.8.14
    diff --git a/docs/functions_func_k.html b/docs/functions_func_k.html new file mode 100644 index 00000000..ced1e9b6 --- /dev/null +++ b/docs/functions_func_k.html @@ -0,0 +1,98 @@ + + + + + + + +Cutlass: Class Members - Functions + + + + + + + + + + +
    +
    +

    @@ -103,7 +103,7 @@ Namespaces

    + + + + + +
    +
    Cutlass +
    +
    CUDA Templates for Linear Algebra Subroutines and Solvers
    +
    + + + + + + + + + + +
    +
    + + +
    + +
    + +
    +  + +

    - k -

    +
    + +
    +Generated on Wed Sep 19 2018 13:58:53 for Cutlass by   +doxygen + 1.8.14 +
    + + diff --git a/docs/functions_func_l.html b/docs/functions_func_l.html index c76f9fc5..76d84a05 100644 --- a/docs/functions_func_l.html +++ b/docs/functions_func_l.html @@ -74,30 +74,56 @@ $(function() {
  • launch() : cutlass::gemm::Gemm< GemmTraits_ >
    • +
  • Launch() +: cutlass::gemm::Launch< Gemm, WithLaunchBounds > +, cutlass::gemm::Launch< Gemm, false > +
  • leading_dim() -: cutlass::TensorRef< Storage_, Rank_ > +: cutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ > +, cutlass::TensorRef< Storage_, Rank_, MapFunc_, 1, Index_, LongIndex_ >
  • LinearScaling() -: cutlass::gemm::LinearScaling< Scalar_, FragmentMultiplyAdd_ > +: cutlass::gemm::LinearScaling< Scalar_, FragmentMultiplyAdd_ > +
    • +
  • LinearScalingDevicePtr() +: cutlass::gemm::LinearScalingDevicePtr< Scalar_, FragmentMultiplyAdd_ >
  • load() -: cutlass::FragmentLoad< IteratorFragment::kScalar, kAccessSize, Scalar_, Memory_, FragmentElement_, kStride > -, cutlass::FragmentLoad< IteratorFragment::kWmmaMatrix, kAccessSize, Scalar_, Memory_, FragmentElement_, kStride > -, cutlass::Load< Scalar_, Lanes_, Memory_, bool, size_t > -, cutlass::Load< double, 2, Memory_, true, 16 > -, cutlass::Load< Scalar_, Lanes_, Memory_, true, 16 > -, cutlass::Load< Scalar_, Lanes_, Memory_, true, 4 > -, cutlass::Load< Scalar_, Lanes_, Memory_, true, 8 > -, cutlass::TileLoadIterator< Traits_, Scalar_, Advance_, MemorySpace, Index_, FragmentElement_, IteratorFragment_, Skew_ > +: cutlass::Load< Scalar_, kAccessSize, Memory_, kFragmentElementType, FragmentElement_, kStride, size > +, cutlass::Load< double, 2, Memory_, FragmentElementType::kScalar, double, kStride, 16 > +, cutlass::Load< Scalar_, kAccessSize, Memory_, FragmentElementType::kScalar, Scalar_, 1, 2 > +, cutlass::Load< Scalar_, kAccessSize, Memory_, FragmentElementType::kScalar, Scalar_, kStride, 16 > +, cutlass::Load< Scalar_, kAccessSize, Memory_, FragmentElementType::kScalar, Scalar_, kStride, 4 > +, cutlass::Load< Scalar_, kAccessSize, Memory_, FragmentElementType::kScalar, Scalar_, kStride, 8 > +, cutlass::Load< Scalar_, kAccessSize, Memory_, FragmentElementType::kWmmaMatrix, FragmentElement_, kStride, size > +, cutlass::Load< Vector< bin1_t, 32 >, kAccessSize, Memory_, FragmentElementType::kWmmaMatrix, FragmentElement_, kStride, size > +, cutlass::Load< Vector< int4_t, 8 >, kAccessSize, Memory_, FragmentElementType::kWmmaMatrix, FragmentElement_, kStride, size > +, cutlass::Load< Vector< uint4_t, 8 >, kAccessSize, Memory_, FragmentElementType::kWmmaMatrix, FragmentElement_, kStride, size > +, cutlass::TileLoadIterator< Traits_, Scalar_, Advance_, MemorySpace, Index_, FragmentElement_, FragmentElementType_, Skew_ > +, cutlass::TileStoreIterator< Traits_, Scalar_, Advance_, MemorySpace, Index_, FragmentElement_, FragmentElementType_, Skew_ > +, cutlass::ZipTileIterator< First_, Second_ > +
    • +
  • load_element() +: cutlass::gemm::GemmGlobalIteratorAb< TileTraits_, Index_ > +, cutlass::gemm::GemmGlobalIteratorCd< TileTraits_, Index_ > +, cutlass::gemm::IgemmGlobalIteratorAb< TileTraits_, Index_ > +, cutlass::gemm::WmmaGemmGlobalIteratorCd< TileTraits_, Index_ > +, cutlass::TileLoadIterator< Traits_, Scalar_, Advance_, MemorySpace, Index_, FragmentElement_, FragmentElementType_, Skew_ > +, cutlass::TileStoreIterator< Traits_, Scalar_, Advance_, MemorySpace, Index_, FragmentElement_, FragmentElementType_, Skew_ >
  • load_post_increment() -: cutlass::TileLoadIterator< Traits_, Scalar_, Advance_, MemorySpace, Index_, FragmentElement_, IteratorFragment_, Skew_ > +: cutlass::gemm::GemmGlobalIteratorAb< TileTraits_, Index_ > +, cutlass::gemm::GemmGlobalIteratorCd< TileTraits_, Index_ > +, cutlass::gemm::WmmaGemmGlobalIteratorCd< TileTraits_, Index_ > +, cutlass::TileLoadIterator< Traits_, Scalar_, Advance_, MemorySpace, Index_, FragmentElement_, FragmentElementType_, Skew_ > +, cutlass::TileStoreIterator< Traits_, Scalar_, Advance_, MemorySpace, Index_, FragmentElement_, FragmentElementType_, Skew_ > +, cutlass::ZipTileIterator< First_, Second_ >
    • -Generated on Thu May 3 2018 16:36:15 for Cutlass by   +Generated on Wed Sep 19 2018 13:58:53 for Cutlass by   doxygen 1.8.14
    diff --git a/docs/functions_func_m.html b/docs/functions_func_m.html index 2c68ec4f..97f28656 100644 --- a/docs/functions_func_m.html +++ b/docs/functions_func_m.html @@ -71,23 +71,38 @@ $(function() {  

    - m -

    -Generated on Thu May 3 2018 16:36:15 for Cutlass by   +Generated on Wed Sep 19 2018 13:58:53 for Cutlass by   doxygen 1.8.14
    diff --git a/docs/functions_func_n.html b/docs/functions_func_n.html new file mode 100644 index 00000000..3b16224a --- /dev/null +++ b/docs/functions_func_n.html @@ -0,0 +1,89 @@ + + + + + + + +Cutlass: Class Members - Functions + + + + + + + + + + +
    +
    + + + + + + +
    +
    Cutlass +
    +
    CUDA Templates for Linear Algebra Subroutines and Solvers
    +
    +
    + + + + + + + +
    + +
    +
    + + +
    + +
    + +
    +  + +

    - n -

    +
    + +
    +Generated on Wed Sep 19 2018 13:58:53 for Cutlass by   +doxygen + 1.8.14 +
    + + diff --git a/docs/functions_func_o.html b/docs/functions_func_o.html index fb7b39f7..a7c0b04d 100644 --- a/docs/functions_func_o.html +++ b/docs/functions_func_o.html @@ -72,8 +72,12 @@ $(function() {

    - o -

    -Generated on Thu May 3 2018 16:36:15 for Cutlass by   +Generated on Wed Sep 19 2018 13:58:53 for Cutlass by   doxygen 1.8.14
    diff --git a/docs/functions_type_f.html b/docs/functions_type_f.html index a71defeb..8ba4932b 100644 --- a/docs/functions_type_f.html +++ b/docs/functions_type_f.html @@ -72,45 +72,59 @@ $(function() {

    - f -

    -Generated on Thu May 3 2018 16:36:15 for Cutlass by   +Generated on Wed Sep 19 2018 13:58:53 for Cutlass by   doxygen 1.8.14
    diff --git a/docs/functions_type_g.html b/docs/functions_type_g.html index 4ae366f3..05c838d3 100644 --- a/docs/functions_type_g.html +++ b/docs/functions_type_g.html @@ -73,19 +73,19 @@ $(function() {

    - g -

    -Generated on Thu May 3 2018 16:36:15 for Cutlass by   +Generated on Wed Sep 19 2018 13:58:53 for Cutlass by   doxygen 1.8.14
    diff --git a/docs/functions_type_k.html b/docs/functions_type_k.html new file mode 100644 index 00000000..b5e5add5 --- /dev/null +++ b/docs/functions_type_k.html @@ -0,0 +1,86 @@ + + + + + + + +Cutlass: Class Members - Typedefs + + + + + + + + + + +
    +
    + + + + + + +
    +
    Cutlass +
    +
    CUDA Templates for Linear Algebra Subroutines and Solvers
    +
    +
    + + + + + + + +
    + +
    +
    + + +
    + +
    + +
    +  + +

    - k -

    +
    + +
    +Generated on Wed Sep 19 2018 13:58:53 for Cutlass by   +doxygen + 1.8.14 +
    + + diff --git a/docs/functions_type_l.html b/docs/functions_type_l.html index 2e7334f0..abda026b 100644 --- a/docs/functions_type_l.html +++ b/docs/functions_type_l.html @@ -72,13 +72,19 @@ $(function() {

    - l -

    -Generated on Thu May 3 2018 16:36:15 for Cutlass by   +Generated on Wed Sep 19 2018 13:58:53 for Cutlass by   doxygen 1.8.14
    diff --git a/docs/functions_type_m.html b/docs/functions_type_m.html index 043340a5..c7639549 100644 --- a/docs/functions_type_m.html +++ b/docs/functions_type_m.html @@ -71,26 +71,33 @@ $(function() {  

    - m -

    -Generated on Thu May 3 2018 16:36:15 for Cutlass by   +Generated on Wed Sep 19 2018 13:58:53 for Cutlass by   doxygen 1.8.14
    diff --git a/docs/functions_type_n.html b/docs/functions_type_n.html index bb5ad36c..187630cf 100644 --- a/docs/functions_type_n.html +++ b/docs/functions_type_n.html @@ -78,7 +78,7 @@ $(function() {
    -Generated on Thu May 3 2018 16:36:15 for Cutlass by   +Generated on Wed Sep 19 2018 13:58:53 for Cutlass by   doxygen 1.8.14
    diff --git a/docs/functions_type_o.html b/docs/functions_type_o.html index 42ed2813..d3d71d55 100644 --- a/docs/functions_type_o.html +++ b/docs/functions_type_o.html @@ -72,7 +72,7 @@ $(function() {

    - o -

    struct  cutlass::Fragment< Element_, kElements_, kAlignment_ >
     A template defining Fragment Concept. More...
     
    struct  cutlass::ZipFragment< First_, Second_ >
     A template defining Fragment Concept. More...
     

    Detailed Description

    Fragment Concept is a statically sized array for storing parts of tiles held by individual CUDA threads.

    @@ -94,7 +97,7 @@ Classes
    - +

    Classes

    struct  cutlass::TileLoadIterator< Traits_, Scalar_, Advance_, MemorySpace, Index_, FragmentElement_, IteratorFragment_, Skew_ >
    struct  cutlass::TileLoadIterator< Traits_, Scalar_, Advance_, MemorySpace, Index_, FragmentElement_, FragmentElementType_, Skew_ >
     An iterator implementing Tile Load Iterator Concept for loading a tile from memory. More...
     
    @@ -96,7 +96,7 @@ Classes
    - +

    Classes

    struct  cutlass::TileStoreIterator< Traits_, Scalar_, Advance_, MemorySpace, Index_, FragmentElement_, IteratorFragment_, Skew_ >
    struct  cutlass::TileStoreIterator< Traits_, Scalar_, Advance_, MemorySpace, Index_, FragmentElement_, FragmentElementType_, Skew_ >
     An iterator implementing Tile Store Iterator Concept for storing a tile to memory. More...
     
    @@ -96,7 +96,7 @@ Classes
    - +

    Classes

    struct  cutlass::TileTraits< Tile_, Delta_, Iterations_, ThreadOffset_ >
    struct  cutlass::TileTraits< Tile_, Delta_, Iterations_, ThreadOffset_, AccessSize >
     A template defining Tile Traits Concept. More...
     
    @@ -93,7 +93,7 @@ Classes
    @@ -107,7 +107,7 @@ Namespaces
    -Generated on Thu May 3 2018 16:36:12 for Cutlass by   +Generated on Wed Sep 19 2018 13:58:50 for Cutlass by   doxygen 1.8.14
    diff --git a/docs/hgemm__global__tile_8h_source.html b/docs/hgemm__global__tile_8h_source.html index bdd647d1..8d7e02f6 100644 --- a/docs/hgemm__global__tile_8h_source.html +++ b/docs/hgemm__global__tile_8h_source.html @@ -76,34 +76,34 @@ $(function() {
    hgemm_global_tile.h
    -Go to the documentation of this file.
    1 /***************************************************************************************************
    2  * Copyright (c) 2017-2018, NVIDIA CORPORATION. All rights reserved.
    3  *
    4  * Redistribution and use in source and binary forms, with or without modification, are permitted
    5  * provided that the following conditions are met:
    6  * * Redistributions of source code must retain the above copyright notice, this list of
    7  * conditions and the following disclaimer.
    8  * * Redistributions in binary form must reproduce the above copyright notice, this list of
    9  * conditions and the following disclaimer in the documentation and/or other materials
    10  * provided with the distribution.
    11  * * Neither the name of the NVIDIA CORPORATION nor the names of its contributors may be used
    12  * to endorse or promote products derived from this software without specific prior written
    13  * permission.
    14  *
    15  * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND ANY EXPRESS OR
    16  * IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND
    17  * FITNESS FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE
    18  * FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING,
    19  * BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS;
    20  * OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT,
    21  * STRICT LIABILITY, OR TOR (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
    22  * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
    23  *
    24  **************************************************************************************************/
    30 #pragma once
    31 
    32 #include <cutlass/coord.h>
    33 #include <cutlass/gemm/gemm_global_tile.h>
    34 #include <cutlass/matrix_traits.h>
    35 #include <cutlass/reshape_tile.h>
    36 
    37 namespace cutlass {
    38 namespace gemm {
    39 
    41 
    42 template <GemmOperand::Kind kOperand_,
    43  MatrixLayout::Kind kLayout_,
    44  typename Scalar_,
    45  typename Tile_,
    46  typename Threads_,
    47  int kAccessSize_>
    48 struct HgemmCrosswiseGlobalTileTraits : public GemmGlobalTileTraits<
    49  // Which GEMM operand?
    50  kOperand_,
    51  // The layout.
    52  kLayout_,
    53  // The scalar.
    54  Scalar_,
    55  // The tile.
    56  Tile_,
    57  // The threads.
    58  Threads_,
    59  // The number of scalars per LDG/STG.
    60  kAccessSize_> {
    62  typedef GemmGlobalTileTraits<kOperand_, kLayout_, Scalar_, Tile_, Threads_, kAccessSize_> Base;
    64  typedef typename Base::Threads Threads;
    66  typedef Shape<1, 2, Base::Tile::kC> ThreadsDelta;
    68  typedef Shape<Base::Threads::kH * 2, 1, Base::Threads::kW, Base::kAccessSize> Delta;
    70  typedef Shape<Base::Tile::kH / Base::Threads::kH / 2,
    71  2,
    72  Base::Tile::kW / Base::Threads::kW,
    73  Base::Tile::kC / Base::kAccessSize>
    74  Iterations;
    76  struct ThreadOffset {
    77  CUTLASS_HOST_DEVICE
    78  Coord<4> operator()() const {
    79  int thread_offset_h = threadIdx.x / Threads::kW * ThreadsDelta::kH;
    80  int thread_offset_w = threadIdx.x % Threads::kW * ThreadsDelta::kW;
    81 
    82  return make_Coord(0, thread_offset_h, thread_offset_w, 0);
    83  }
    84  };
    85 };
    86 
    88 
    89 } // namespace gemm
    90 } // namespace cutlass
    cutlass
    Definition: convert.h:33
    +Go to the documentation of this file.
    1 /***************************************************************************************************
    2  * Copyright (c) 2017-2018, NVIDIA CORPORATION. All rights reserved.
    3  *
    4  * Redistribution and use in source and binary forms, with or without modification, are permitted
    5  * provided that the following conditions are met:
    6  * * Redistributions of source code must retain the above copyright notice, this list of
    7  * conditions and the following disclaimer.
    8  * * Redistributions in binary form must reproduce the above copyright notice, this list of
    9  * conditions and the following disclaimer in the documentation and/or other materials
    10  * provided with the distribution.
    11  * * Neither the name of the NVIDIA CORPORATION nor the names of its contributors may be used
    12  * to endorse or promote products derived from this software without specific prior written
    13  * permission.
    14  *
    15  * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND ANY EXPRESS OR
    16  * IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND
    17  * FITNESS FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE
    18  * FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING,
    19  * BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS;
    20  * OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT,
    21  * STRICT LIABILITY, OR TOR (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
    22  * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
    23  *
    24  **************************************************************************************************/
    30 #pragma once
    31 
    32 #include "cutlass/coord.h"
    33 #include "cutlass/gemm/gemm_global_tile.h"
    34 #include "cutlass/matrix_traits.h"
    35 #include "cutlass/reshape_tile.h"
    36 
    37 namespace cutlass {
    38 namespace gemm {
    39 
    41 
    42 template <GemmOperand::Kind kOperand_,
    43  MatrixLayout::Kind kLayout_,
    44  typename Scalar_,
    45  typename Tile_,
    46  typename Threads_,
    47  int kAccessSize_>
    48 struct HgemmCrosswiseGlobalTileTraits : public GemmGlobalTileTraits<
    49  // Which GEMM operand?
    50  kOperand_,
    51  // The layout.
    52  kLayout_,
    53  // The scalar.
    54  Scalar_,
    55  // The tile.
    56  Tile_,
    57  // The threads.
    58  Threads_,
    59  // The number of scalars per LDG/STG.
    60  kAccessSize_> {
    62  typedef GemmGlobalTileTraits<kOperand_, kLayout_, Scalar_, Tile_, Threads_, kAccessSize_> Base;
    64  typedef typename Base::Threads Threads;
    66  typedef Shape<1, 2, Base::VectorizedTile::kC> ThreadsDelta;
    68  typedef Shape<Base::Threads::kH * 2, 1, Base::Threads::kW, Base::kAccessSize> Delta;
    70  typedef Shape<Base::VectorizedTile::kH / Base::Threads::kH / 2,
    71  2,
    72  Base::VectorizedTile::kW / Base::Threads::kW,
    73  Base::VectorizedTile::kC / Base::kAccessSize>
    74  Iterations;
    76  struct ThreadOffset {
    77  CUTLASS_HOST_DEVICE
    78  Coord<4> operator()() const {
    79  int thread_offset_h = threadIdx.x / Threads::kW * ThreadsDelta::kH;
    80  int thread_offset_w = threadIdx.x % Threads::kW * ThreadsDelta::kW;
    81 
    82  return make_Coord(0, thread_offset_h, thread_offset_w, 0);
    83  }
    84  };
    85 };
    86 
    88 
    89 } // namespace gemm
    90 } // namespace cutlass
    cutlass::gemm::HgemmCrosswiseGlobalTileTraits::Iterations
    Shape< Base::VectorizedTile::kH/Base::Threads::kH/2, 2, Base::VectorizedTile::kW/Base::Threads::kW, Base::VectorizedTile::kC/Base::kAccessSize > Iterations
    The number of iterations needed to load/store the tile.
    Definition: hgemm_global_tile.h:74
    +
    cutlass
    Definition: convert.h:33
    gemm_global_tile.h
    Defines iterators for efficiently loading and storing to global memory.
    cutlass::gemm::GemmGlobalTileTraits
    Definition: gemm_global_tile.h:70
    coord.h
    A Coord is a coordinate of arbitrary rank into a tensor or matrix.
    -
    cutlass::make_Coord
    CUTLASS_HOST_DEVICE Coord< 1 > make_Coord(int _0)
    Helper to make a 2-element coordinate.
    Definition: coord.h:241
    -
    cutlass::gemm::HgemmCrosswiseGlobalTileTraits::Iterations
    Shape< Base::Tile::kH/Base::Threads::kH/2, 2, Base::Tile::kW/Base::Threads::kW, Base::Tile::kC/Base::kAccessSize > Iterations
    The number of iterations needed to load/store the tile.
    Definition: hgemm_global_tile.h:74
    +
    cutlass::make_Coord
    CUTLASS_HOST_DEVICE Coord< 1 > make_Coord(int _0)
    Helper to make a 2-element coordinate.
    Definition: coord.h:318
    +
    cutlass::gemm::HgemmCrosswiseGlobalTileTraits::ThreadsDelta
    Shape< 1, 2, Base::VectorizedTile::kC > ThreadsDelta
    The threads strides.
    Definition: hgemm_global_tile.h:66
    cutlass::gemm::HgemmCrosswiseGlobalTileTraits::Threads
    Base::Threads Threads
    The threads.
    Definition: hgemm_global_tile.h:64
    cutlass::Shape::kH
    static int const kH
    The height of the cube.
    Definition: shape.h:68
    cutlass::gemm::HgemmCrosswiseGlobalTileTraits::ThreadOffset::operator()
    CUTLASS_HOST_DEVICE Coord< 4 > operator()() const
    Definition: hgemm_global_tile.h:78
    cutlass::gemm::HgemmCrosswiseGlobalTileTraits::Delta
    Shape< Base::Threads::kH *2, 1, Base::Threads::kW, Base::kAccessSize > Delta
    The strides in each dimension between different loads/stores.
    Definition: hgemm_global_tile.h:68
    -
    cutlass::gemm::HgemmCrosswiseGlobalTileTraits::ThreadsDelta
    Shape< 1, 2, Base::Tile::kC > ThreadsDelta
    The threads strides.
    Definition: hgemm_global_tile.h:66
    reshape_tile.h
    Defines a type for restructuring a tile.
    CUTLASS_HOST_DEVICE
    #define CUTLASS_HOST_DEVICE
    Definition: cutlass.h:46
    cutlass::gemm::HgemmCrosswiseGlobalTileTraits::Base
    GemmGlobalTileTraits< kOperand_, kLayout_, Scalar_, Tile_, Threads_, kAccessSize_ > Base
    The base class.
    Definition: hgemm_global_tile.h:62
    cutlass::gemm::HgemmCrosswiseGlobalTileTraits
    Definition: hgemm_global_tile.h:48
    cutlass::Shape
    A Shape implementing Layout Concept describing the dimensions of a cube.
    Definition: shape.h:64
    +
    cutlass::gemm::GemmGlobalTileTraits::Threads
    ReshapeThreads< VectorizedTile, Threads_ >::Threads Threads
    The threads shape.
    Definition: gemm_global_tile.h:88
    cutlass::Shape::kW
    static int const kW
    The width of the cube.
    Definition: shape.h:70
    -
    cutlass::MatrixLayout::Kind
    Kind
    Definition: matrix_traits.h:36
    +
    cutlass::MatrixLayout::Kind
    Kind
    Enumeration defining fundamental contiguous layouts.
    Definition: matrix_traits.h:159
    cutlass::gemm::GemmGlobalTileTraits::kAccessSize
    static int const kAccessSize
    The number of scalars per LDG/STG.
    Definition: gemm_global_tile.h:80
    cutlass::gemm::HgemmCrosswiseGlobalTileTraits::ThreadOffset
    Computes the thread offset in (H, W) based on thread ID.
    Definition: hgemm_global_tile.h:76
    -
    cutlass::GemmOperand::Kind
    Kind
    Definition: matrix_traits.h:43
    -
    cutlass::gemm::GemmGlobalTileTraits::Threads
    ReshapeThreads< Tile, Threads_ >::Threads Threads
    The threads shape.
    Definition: gemm_global_tile.h:87
    +
    cutlass::GemmOperand::Kind
    Kind
    Definition: matrix_traits.h:357
    matrix_traits.h
    Defines properties of matrices used to denote layout and operands to GEMM kernels.
    -Generated on Thu May 3 2018 16:36:12 for Cutlass by   +Generated on Wed Sep 19 2018 13:58:50 for Cutlass by   doxygen 1.8.14
    diff --git a/docs/hgemm__multiply__add_8h.html b/docs/hgemm__multiply__add_8h.html index 3c6c609e..41ba8db9 100644 --- a/docs/hgemm__multiply__add_8h.html +++ b/docs/hgemm__multiply__add_8h.html @@ -82,15 +82,15 @@ $(function() {

    Specialization implementing multiply-add operation on half-precision floating point fragments. More...

    -
    #include <cutlass/fragment.h>
    -#include <cutlass/gemm/thread_multiply_add.h>
    +
    #include "cutlass/fragment.h"
    +#include "cutlass/gemm/thread_multiply_add.h"

    Go to the source code of this file.

    - - + +

    Classes

    struct  cutlass::gemm::ThreadMultiplyAdd< AccumulatorsPerThread_, ThreadsPerWarp_, half, half, half >
     Template performing matrix multiply-add operation within a thread. More...
    struct  cutlass::gemm::ThreadMultiplyAdd< ThreadGemmShape_, ThreadsPerWarp_, half, half, half >
     Template performing matrix multiply-add operation within a thread. More...
     
    -Generated on Thu May 3 2018 16:36:12 for Cutlass by   +Generated on Wed Sep 19 2018 13:58:50 for Cutlass by   doxygen 1.8.14
    diff --git a/docs/hgemm__multiply__add_8h_source.html b/docs/hgemm__multiply__add_8h_source.html index 73ef9040..40e849bf 100644 --- a/docs/hgemm__multiply__add_8h_source.html +++ b/docs/hgemm__multiply__add_8h_source.html @@ -76,30 +76,31 @@ $(function() {
    hgemm_multiply_add.h
    -Go to the documentation of this file.
    1 /***************************************************************************************************
    2  * Copyright (c) 2017-2018, NVIDIA CORPORATION. All rights reserved.
    3  *
    4  * Redistribution and use in source and binary forms, with or without modification, are permitted
    5  * provided that the following conditions are met:
    6  * * Redistributions of source code must retain the above copyright notice, this list of
    7  * conditions and the following disclaimer.
    8  * * Redistributions in binary form must reproduce the above copyright notice, this list of
    9  * conditions and the following disclaimer in the documentation and/or other materials
    10  * provided with the distribution.
    11  * * Neither the name of the NVIDIA CORPORATION nor the names of its contributors may be used
    12  * to endorse or promote products derived from this software without specific prior written
    13  * permission.
    14  *
    15  * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND ANY EXPRESS OR
    16  * IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND
    17  * FITNESS FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE
    18  * FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING,
    19  * BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS;
    20  * OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT,
    21  * STRICT LIABILITY, OR TOR (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
    22  * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
    23  *
    24  **************************************************************************************************/
    29 #pragma once
    30 
    31 #include <cutlass/fragment.h>
    32 
    33 #include <cutlass/gemm/thread_multiply_add.h>
    34 
    35 namespace cutlass {
    36 namespace gemm {
    37 
    39 
    41 template <typename AccumulatorsPerThread_, typename ThreadsPerWarp_>
    42 struct ThreadMultiplyAdd<AccumulatorsPerThread_, ThreadsPerWarp_, half, half, half> {
    44  typedef Shape<1, 1, 2, 1> InstructionShape;
    46  typedef AccumulatorsPerThread_ AccumulatorsPerThread;
    48  typedef ThreadsPerWarp_ ThreadsPerWarp;
    50  typedef typename ShapeMul<AccumulatorsPerThread, ThreadsPerWarp>::Shape AccumulatorsPerWarp;
    52  typedef half ScalarA;
    54  typedef Fragment<ScalarA, AccumulatorsPerThread::kW> FragmentA;
    56  typedef half ScalarB;
    58  typedef Fragment<ScalarB, AccumulatorsPerThread::kH> FragmentB;
    60  typedef half ScalarC;
    62  typedef Fragment<half, AccumulatorsPerThread::kH * AccumulatorsPerThread::kW> Accumulators;
    63 
    65  static_assert(AccumulatorsPerThread::kH % 2 == 0, "Invalid size");
    66  static_assert(AccumulatorsPerThread::kW % 2 == 0, "Invalid size");
    67 
    69  CUTLASS_DEVICE ThreadMultiplyAdd() {}
    70 
    72  CUTLASS_DEVICE void multiply_add(FragmentA const& a,
    73  FragmentB const& b,
    74  Accumulators const& c,
    75  Accumulators& d) {
    76 #if defined(__CUDACC__) && __CUDA_ARCH__ >= 530
    77  // The inputs.
    78  __half2 const* a_half2 = reinterpret_cast<__half2 const*>(&a[0]);
    79  __half2 const* b_half2 = reinterpret_cast<__half2 const*>(&b[0]);
    80  __half2 const* c_half2 = reinterpret_cast<__half2 const*>(&c[0]);
    81 
    82  // The output.
    83  __half2* d_half2 = reinterpret_cast<__half2*>(&d[0]);
    84 
    85  for (int j = 0; j < AccumulatorsPerThread::kH / 2; ++j) {
    86  for (int i = 0; i < AccumulatorsPerThread::kW / 2; ++i) {
    87  // The offsets in the output fragment.
    88  int const k0 = (2 * j + 0) * (AccumulatorsPerThread::kW / 2) + i;
    89  int const k1 = (2 * j + 1) * (AccumulatorsPerThread::kW / 2) + i;
    90 
    91  // Compute the product a[i] * b[j].H0_H0.
    92  d_half2[k0] = __hfma2(a_half2[i], __low2half2(b_half2[j]), c_half2[k0]);
    93  // Compute the product a[i] * b[j].H1_H1.
    94  d_half2[k1] = __hfma2(a_half2[i], __high2half2(b_half2[j]), c_half2[k1]);
    95  }
    96  }
    97 #endif
    98  }
    99 };
    100 
    102 
    103 } // namespace gemm
    104 } // namespace cutlass
    cutlass::gemm::ThreadMultiplyAdd< AccumulatorsPerThread_, ThreadsPerWarp_, half, half, half >::ScalarA
    half ScalarA
    The type for A.
    Definition: hgemm_multiply_add.h:52
    +Go to the documentation of this file.
    1 /***************************************************************************************************
    2  * Copyright (c) 2017-2018, NVIDIA CORPORATION. All rights reserved.
    3  *
    4  * Redistribution and use in source and binary forms, with or without modification, are permitted
    5  * provided that the following conditions are met:
    6  * * Redistributions of source code must retain the above copyright notice, this list of
    7  * conditions and the following disclaimer.
    8  * * Redistributions in binary form must reproduce the above copyright notice, this list of
    9  * conditions and the following disclaimer in the documentation and/or other materials
    10  * provided with the distribution.
    11  * * Neither the name of the NVIDIA CORPORATION nor the names of its contributors may be used
    12  * to endorse or promote products derived from this software without specific prior written
    13  * permission.
    14  *
    15  * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND ANY EXPRESS OR
    16  * IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND
    17  * FITNESS FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE
    18  * FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING,
    19  * BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS;
    20  * OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT,
    21  * STRICT LIABILITY, OR TOR (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
    22  * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
    23  *
    24  **************************************************************************************************/
    29 #pragma once
    30 
    31 #include "cutlass/fragment.h"
    32 
    33 #include "cutlass/gemm/thread_multiply_add.h"
    34 
    35 namespace cutlass {
    36 namespace gemm {
    37 
    39 
    41 template <typename ThreadGemmShape_, typename ThreadsPerWarp_>
    42 struct ThreadMultiplyAdd<ThreadGemmShape_, ThreadsPerWarp_, half, half, half> {
    44  typedef Shape<1, 1, 2, 1> InstructionShape;
    46  typedef ThreadGemmShape_ ThreadGemmShape;
    48  typedef ThreadGemmShape AccumulatorsPerThread;
    50  typedef ThreadsPerWarp_ ThreadsPerWarp;
    52  typedef typename ShapeMul<ThreadGemmShape, ThreadsPerWarp>::Shape AccumulatorsPerWarp;
    54  typedef half ScalarA;
    56  typedef Fragment<ScalarA, AccumulatorsPerThread::kW> FragmentA;
    58  typedef half ScalarB;
    60  typedef Fragment<ScalarB, AccumulatorsPerThread::kH> FragmentB;
    62  typedef half ScalarC;
    64  typedef Fragment<half, AccumulatorsPerThread::kH * AccumulatorsPerThread::kW> Accumulators;
    65 
    67  static_assert(AccumulatorsPerThread::kH % 2 == 0, "Invalid size");
    68  static_assert(AccumulatorsPerThread::kW % 2 == 0, "Invalid size");
    69 
    71  CUTLASS_DEVICE ThreadMultiplyAdd() {}
    72 
    74  CUTLASS_DEVICE void multiply_add(FragmentA const& a,
    75  FragmentB const& b,
    76  Accumulators const& c,
    77  Accumulators& d) {
    78 #if defined(__CUDACC__) && __CUDA_ARCH__ >= 530
    79  // The inputs.
    80  __half2 const* a_half2 = reinterpret_cast<__half2 const*>(&a[0]);
    81  __half2 const* b_half2 = reinterpret_cast<__half2 const*>(&b[0]);
    82  __half2 const* c_half2 = reinterpret_cast<__half2 const*>(&c[0]);
    83 
    84  // The output.
    85  __half2* d_half2 = reinterpret_cast<__half2*>(&d[0]);
    86 
    87  for (int j = 0; j < AccumulatorsPerThread::kH / 2; ++j) {
    88  for (int i = 0; i < AccumulatorsPerThread::kW / 2; ++i) {
    89  // The offsets in the output fragment.
    90  int const k0 = (2 * j + 0) * (AccumulatorsPerThread::kW / 2) + i;
    91  int const k1 = (2 * j + 1) * (AccumulatorsPerThread::kW / 2) + i;
    92 
    93  // Compute the product a[i] * b[j].low.
    94  d_half2[k0] = __hfma2(a_half2[i], __low2half2(b_half2[j]), c_half2[k0]);
    95  // Compute the product a[i] * b[j].high.
    96  d_half2[k1] = __hfma2(a_half2[i], __high2half2(b_half2[j]), c_half2[k1]);
    97  }
    98  }
    99 #endif
    100  }
    101 };
    102 
    104 
    105 } // namespace gemm
    106 } // namespace cutlass
    cutlass::gemm::ThreadMultiplyAdd< ThreadGemmShape_, ThreadsPerWarp_, half, half, half >::ThreadMultiplyAdd
    CUTLASS_DEVICE ThreadMultiplyAdd()
    Make sure there&#39;s an even number of elements in both dimensions.
    Definition: hgemm_multiply_add.h:71
    +
    cutlass::gemm::ThreadMultiplyAdd< ThreadGemmShape_, ThreadsPerWarp_, half, half, half >::ScalarC
    half ScalarC
    The type for C and D.
    Definition: hgemm_multiply_add.h:62
    cutlass
    Definition: convert.h:33
    -
    cutlass::gemm::ThreadMultiplyAdd< AccumulatorsPerThread_, ThreadsPerWarp_, half, half, half >::Accumulators
    Fragment< half, AccumulatorsPerThread::kH *AccumulatorsPerThread::kW > Accumulators
    The accumulators.
    Definition: hgemm_multiply_add.h:62
    -
    cutlass::gemm::ThreadMultiplyAdd< AccumulatorsPerThread_, ThreadsPerWarp_, half, half, half >::AccumulatorsPerWarp
    ShapeMul< AccumulatorsPerThread, ThreadsPerWarp >::Shape AccumulatorsPerWarp
    The number of accumulators per warp.
    Definition: hgemm_multiply_add.h:50
    -
    cutlass::gemm::ThreadMultiplyAdd< AccumulatorsPerThread_, ThreadsPerWarp_, half, half, half >::ScalarC
    half ScalarC
    The type for C and D.
    Definition: hgemm_multiply_add.h:60
    -
    cutlass::gemm::ThreadMultiplyAdd< AccumulatorsPerThread_, ThreadsPerWarp_, half, half, half >::ThreadMultiplyAdd
    CUTLASS_DEVICE ThreadMultiplyAdd()
    Make sure there&#39;s an even number of elements in both dimensions.
    Definition: hgemm_multiply_add.h:69
    +
    cutlass::gemm::ThreadMultiplyAdd< ThreadGemmShape_, ThreadsPerWarp_, half, half, half >::FragmentB
    Fragment< ScalarB, AccumulatorsPerThread::kH > FragmentB
    The fragment for B.
    Definition: hgemm_multiply_add.h:60
    +
    cutlass::gemm::ThreadMultiplyAdd< ThreadGemmShape_, ThreadsPerWarp_, half, half, half >::ThreadGemmShape
    ThreadGemmShape_ ThreadGemmShape
    The number of accumulators per thread.
    Definition: hgemm_multiply_add.h:46
    +
    cutlass::gemm::ThreadMultiplyAdd< ThreadGemmShape_, ThreadsPerWarp_, half, half, half >::ScalarB
    half ScalarB
    The type for B.
    Definition: hgemm_multiply_add.h:58
    cutlass::ShapeMul::Shape
    Shape< A_::kD *B_::kD, A_::kH *B_::kH, A_::kW *B_::kW, A_::kC *B_::kC > Shape
    Definition: shape.h:119
    cutlass::Fragment
    A template defining Fragment Concept.
    Definition: fragment.h:99
    thread_multiply_add.h
    Template implementing matrix multiply-add operations on fragments.
    -
    cutlass::gemm::ThreadMultiplyAdd< AccumulatorsPerThread_, ThreadsPerWarp_, half, half, half >::InstructionShape
    Shape< 1, 1, 2, 1 > InstructionShape
    The shape of the instruction.
    Definition: hgemm_multiply_add.h:44
    -
    cutlass::gemm::ThreadMultiplyAdd< AccumulatorsPerThread_, ThreadsPerWarp_, half, half, half >::ScalarB
    half ScalarB
    The type for B.
    Definition: hgemm_multiply_add.h:56
    -
    cutlass::gemm::ThreadMultiplyAdd< AccumulatorsPerThread_, ThreadsPerWarp_, half, half, half >::ThreadsPerWarp
    ThreadsPerWarp_ ThreadsPerWarp
    The number of threads per warp.
    Definition: hgemm_multiply_add.h:48
    -
    cutlass::gemm::ThreadMultiplyAdd< AccumulatorsPerThread_, ThreadsPerWarp_, half, half, half >::AccumulatorsPerThread
    AccumulatorsPerThread_ AccumulatorsPerThread
    The number of accumulators per thread.
    Definition: hgemm_multiply_add.h:46
    -
    static_assert
    #define static_assert(__e, __m)
    Definition: platform.h:145
    -
    cutlass::gemm::ThreadMultiplyAdd< AccumulatorsPerThread_, ThreadsPerWarp_, half, half, half >::multiply_add
    CUTLASS_DEVICE void multiply_add(FragmentA const &a, FragmentB const &b, Accumulators const &c, Accumulators &d)
    Multiply : d = a*b + c.
    Definition: hgemm_multiply_add.h:72
    +
    cutlass::gemm::ThreadMultiplyAdd< ThreadGemmShape_, ThreadsPerWarp_, half, half, half >::InstructionShape
    Shape< 1, 1, 2, 1 > InstructionShape
    The shape of the instruction.
    Definition: hgemm_multiply_add.h:44
    +
    cutlass::gemm::ThreadMultiplyAdd< ThreadGemmShape_, ThreadsPerWarp_, half, half, half >::AccumulatorsPerWarp
    ShapeMul< ThreadGemmShape, ThreadsPerWarp >::Shape AccumulatorsPerWarp
    The number of accumulators per warp.
    Definition: hgemm_multiply_add.h:52
    +
    cutlass::gemm::ThreadMultiplyAdd< ThreadGemmShape_, ThreadsPerWarp_, half, half, half >::FragmentA
    Fragment< ScalarA, AccumulatorsPerThread::kW > FragmentA
    The fragment for A.
    Definition: hgemm_multiply_add.h:56
    +
    cutlass::gemm::ThreadMultiplyAdd< ThreadGemmShape_, ThreadsPerWarp_, half, half, half >::Accumulators
    Fragment< half, AccumulatorsPerThread::kH *AccumulatorsPerThread::kW > Accumulators
    The accumulators.
    Definition: hgemm_multiply_add.h:64
    +
    cutlass::gemm::ThreadMultiplyAdd< ThreadGemmShape_, ThreadsPerWarp_, half, half, half >::multiply_add
    CUTLASS_DEVICE void multiply_add(FragmentA const &a, FragmentB const &b, Accumulators const &c, Accumulators &d)
    Multiply : d = a*b + c.
    Definition: hgemm_multiply_add.h:74
    +
    cutlass::gemm::ThreadMultiplyAdd< ThreadGemmShape_, ThreadsPerWarp_, half, half, half >::ScalarA
    half ScalarA
    The type for A.
    Definition: hgemm_multiply_add.h:54
    +
    static_assert
    #define static_assert(__e, __m)
    Definition: platform.h:153
    cutlass::Shape
    A Shape implementing Layout Concept describing the dimensions of a cube.
    Definition: shape.h:64
    -
    cutlass::gemm::ThreadMultiplyAdd
    Template performing matrix multiply-add operation within a thread.
    Definition: thread_multiply_add.h:43
    +
    cutlass::gemm::ThreadMultiplyAdd
    Template performing matrix multiply-add operation within a thread.
    Definition: thread_multiply_add.h:44
    +
    cutlass::gemm::ThreadMultiplyAdd< ThreadGemmShape_, ThreadsPerWarp_, half, half, half >::AccumulatorsPerThread
    ThreadGemmShape AccumulatorsPerThread
    Aliased for compatibility. Will be removed for CUTLASS v2.0.
    Definition: hgemm_multiply_add.h:48
    +
    cutlass::gemm::ThreadMultiplyAdd< ThreadGemmShape_, ThreadsPerWarp_, half, half, half >::ThreadsPerWarp
    ThreadsPerWarp_ ThreadsPerWarp
    The number of threads per warp.
    Definition: hgemm_multiply_add.h:50
    fragment.h
    Defines Fragment, a statically-sized array for storing parts of matrices within a thread&#39;s registers...
    -
    cutlass::gemm::ThreadMultiplyAdd< AccumulatorsPerThread_, ThreadsPerWarp_, half, half, half >::FragmentA
    Fragment< ScalarA, AccumulatorsPerThread::kW > FragmentA
    The fragment for A.
    Definition: hgemm_multiply_add.h:54
    -
    cutlass::gemm::ThreadMultiplyAdd< AccumulatorsPerThread_, ThreadsPerWarp_, half, half, half >::FragmentB
    Fragment< ScalarB, AccumulatorsPerThread::kH > FragmentB
    The fragment for B.
    Definition: hgemm_multiply_add.h:58
    -Generated on Thu May 3 2018 16:36:12 for Cutlass by   +Generated on Wed Sep 19 2018 13:58:50 for Cutlass by   doxygen 1.8.14
    diff --git a/docs/hgemm__swizzle_8h.html b/docs/hgemm__swizzle_8h.html index aef7ac75..93938799 100644 --- a/docs/hgemm__swizzle_8h.html +++ b/docs/hgemm__swizzle_8h.html @@ -83,7 +83,7 @@ $(function() {

    Transposes a tile of 16b elements. Used by HGEMM to construct a K-strided layout in shared memory for multiplicands. More...

    #include <cuda_fp16.h>
    -#include <cutlass/fragment.h>
    +#include "cutlass/fragment.h"

    Go to the source code of this file.

    @@ -103,7 +103,7 @@ Namespaces

    @@ -102,7 +102,7 @@ Namespaces
    -Generated on Thu May 3 2018 16:36:12 for Cutlass by   +Generated on Wed Sep 19 2018 13:58:50 for Cutlass by   doxygen 1.8.14
    diff --git a/docs/hgemm__swizzle_8h_source.html b/docs/hgemm__swizzle_8h_source.html index bb76b510..d882c10f 100644 --- a/docs/hgemm__swizzle_8h_source.html +++ b/docs/hgemm__swizzle_8h_source.html @@ -76,14 +76,14 @@ $(function() {
    hgemm_swizzle.h
    -Go to the documentation of this file.
    1 /***************************************************************************************************
    2  * Copyright (c) 2017-2018, NVIDIA CORPORATION. All rights reserved.
    3  *
    4  * Redistribution and use in source and binary forms, with or without modification, are permitted
    5  * provided that the following conditions are met:
    6  * * Redistributions of source code must retain the above copyright notice, this list of
    7  * conditions and the following disclaimer.
    8  * * Redistributions in binary form must reproduce the above copyright notice, this list of
    9  * conditions and the following disclaimer in the documentation and/or other materials
    10  * provided with the distribution.
    11  * * Neither the name of the NVIDIA CORPORATION nor the names of its contributors may be used
    12  * to endorse or promote products derived from this software without specific prior written
    13  * permission.
    14  *
    15  * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND ANY EXPRESS OR
    16  * IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND
    17  * FITNESS FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE
    18  * FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING,
    19  * BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS;
    20  * OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT,
    21  * STRICT LIABILITY, OR TOR (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
    22  * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
    23  *
    24  **************************************************************************************************/
    29 #pragma once
    30 
    31 #include <cuda_fp16.h>
    32 #include <cutlass/fragment.h>
    33 
    34 namespace cutlass {
    35 namespace gemm {
    36 
    38 
    39 template <typename GlobalIterator_>
    40 struct HgemmSwizzle {
    42  typedef GlobalIterator_ GlobalIterator;
    44  typedef typename GlobalIterator::Fragment Fragment;
    46  typedef typename GlobalIterator::FragmentShape FragmentShape;
    47 
    49  typedef Fragment InputFragment;
    51  typedef Fragment OutputFragment;
    52 
    54  static_assert((platform::is_same<typename Fragment::Element, half>::value), "Works on half");
    55 
    57  static_assert(FragmentShape::kH == 2 && ShapeCount<FragmentShape>::kWc == 2, "Not multiple of 2");
    58 
    60  CUTLASS_DEVICE HgemmSwizzle() {}
    61 
    63  CUTLASS_DEVICE void transform(Fragment const& src, Fragment& dst) {
    64  // Expose src/dst as int arrays.
    65  int const* src_int = reinterpret_cast<int const*>(&src[0]);
    66  int* dst_int = reinterpret_cast<int*>(&dst[0]);
    67 
    68  // Transpose the data.
    69  for (int d = 0; d < FragmentShape::kD; ++d) {
    70  // The indices to read two consecutive "rows".
    71  int const i0 = 2 * d + 0;
    72  int const i1 = 2 * d + 1;
    73 
    74  int a0 = src_int[i0];
    75  int a1 = src_int[i1];
    76 
    77  int b0, b1;
    78  asm volatile("prmt.b32 %0, %1, %2, 0x5410;" : "=r"(b0) : "r"(a0), "r"(a1));
    79  asm volatile("prmt.b32 %0, %1, %2, 0x7632;" : "=r"(b1) : "r"(a0), "r"(a1));
    80 
    81  // The indices to store with "strides".
    82  int const j0 = 0 * (ShapeCount<FragmentShape>::kDhw / 2) + d;
    83  int const j1 = 1 * (ShapeCount<FragmentShape>::kDhw / 2) + d;
    84 
    85  dst_int[j0] = b0;
    86  dst_int[j1] = b1;
    87  }
    88  }
    89 };
    90 
    92 
    93 } // namespace gemm
    94 } // namespace cutlass
    cutlass::gemm::HgemmSwizzle::GlobalIterator
    GlobalIterator_ GlobalIterator
    The global iterator.
    Definition: hgemm_swizzle.h:42
    +Go to the documentation of this file.
    1 /***************************************************************************************************
    2  * Copyright (c) 2017-2018, NVIDIA CORPORATION. All rights reserved.
    3  *
    4  * Redistribution and use in source and binary forms, with or without modification, are permitted
    5  * provided that the following conditions are met:
    6  * * Redistributions of source code must retain the above copyright notice, this list of
    7  * conditions and the following disclaimer.
    8  * * Redistributions in binary form must reproduce the above copyright notice, this list of
    9  * conditions and the following disclaimer in the documentation and/or other materials
    10  * provided with the distribution.
    11  * * Neither the name of the NVIDIA CORPORATION nor the names of its contributors may be used
    12  * to endorse or promote products derived from this software without specific prior written
    13  * permission.
    14  *
    15  * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND ANY EXPRESS OR
    16  * IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND
    17  * FITNESS FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE
    18  * FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING,
    19  * BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS;
    20  * OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT,
    21  * STRICT LIABILITY, OR TOR (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
    22  * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
    23  *
    24  **************************************************************************************************/
    29 #pragma once
    30 
    31 #include <cuda_fp16.h>
    32 #include "cutlass/fragment.h"
    33 
    34 namespace cutlass {
    35 namespace gemm {
    36 
    38 
    39 template <typename GlobalIterator_>
    40 struct HgemmSwizzle {
    42  typedef GlobalIterator_ GlobalIterator;
    44  typedef typename GlobalIterator::Fragment Fragment;
    46  typedef typename GlobalIterator::FragmentShape FragmentShape;
    47 
    49  typedef Fragment InputFragment;
    51  typedef Fragment OutputFragment;
    52 
    54  static_assert((platform::is_same<typename Fragment::Element, half>::value), "Works on half");
    55 
    57  static_assert(FragmentShape::kH == 2 && ShapeCount<FragmentShape>::kWc == 2, "Not multiple of 2");
    58 
    60  CUTLASS_DEVICE HgemmSwizzle() {}
    61 
    63  CUTLASS_DEVICE void transform(Fragment const& src, Fragment& dst) {
    64  // Expose src/dst as int arrays.
    65  int const* src_int = reinterpret_cast<int const*>(&src[0]);
    66  int* dst_int = reinterpret_cast<int*>(&dst[0]);
    67 
    68  // Transpose the data.
    69  for (int d = 0; d < FragmentShape::kD; ++d) {
    70  // The indices to read two consecutive "rows".
    71  int const i0 = 2 * d + 0;
    72  int const i1 = 2 * d + 1;
    73 
    74  int a0 = src_int[i0];
    75  int a1 = src_int[i1];
    76 
    77  int b0, b1;
    78  asm volatile("prmt.b32 %0, %1, %2, 0x5410;" : "=r"(b0) : "r"(a0), "r"(a1));
    79  asm volatile("prmt.b32 %0, %1, %2, 0x7632;" : "=r"(b1) : "r"(a0), "r"(a1));
    80 
    81  // The indices to store with "strides".
    82  int const j0 = 0 * (ShapeCount<FragmentShape>::kDhw / 2) + d;
    83  int const j1 = 1 * (ShapeCount<FragmentShape>::kDhw / 2) + d;
    84 
    85  dst_int[j0] = b0;
    86  dst_int[j1] = b1;
    87  }
    88  }
    89 };
    90 
    92 
    93 } // namespace gemm
    94 } // namespace cutlass
    cutlass::gemm::HgemmSwizzle::GlobalIterator
    GlobalIterator_ GlobalIterator
    The global iterator.
    Definition: hgemm_swizzle.h:42
    cutlass
    Definition: convert.h:33
    -
    cutlass::platform::is_same
    std::is_same (false specialization)
    Definition: platform.h:412
    +
    cutlass::platform::is_same
    std::is_same (false specialization)
    Definition: platform.h:420
    cutlass::gemm::HgemmSwizzle::HgemmSwizzle
    CUTLASS_DEVICE HgemmSwizzle()
    The src/dst must be half fragments.
    Definition: hgemm_swizzle.h:60
    cutlass::gemm::HgemmSwizzle::transform
    CUTLASS_DEVICE void transform(Fragment const &src, Fragment &dst)
    Transform a fragment.
    Definition: hgemm_swizzle.h:63
    cutlass::gemm::HgemmSwizzle::InputFragment
    Fragment InputFragment
    The input fragment.
    Definition: hgemm_swizzle.h:49
    cutlass::gemm::HgemmSwizzle::OutputFragment
    Fragment OutputFragment
    The output fragment.
    Definition: hgemm_swizzle.h:51
    -
    static_assert
    #define static_assert(__e, __m)
    Definition: platform.h:145
    +
    static_assert
    #define static_assert(__e, __m)
    Definition: platform.h:153
    cutlass::gemm::HgemmSwizzle::Fragment
    GlobalIterator::Fragment Fragment
    The source fragment.
    Definition: hgemm_swizzle.h:44
    fragment.h
    Defines Fragment, a statically-sized array for storing parts of matrices within a thread&#39;s registers...
    cutlass::gemm::HgemmSwizzle::FragmentShape
    GlobalIterator::FragmentShape FragmentShape
    The shape of the source fragment.
    Definition: hgemm_swizzle.h:46
    @@ -92,7 +92,7 @@ $(function() {
    -Generated on Thu May 3 2018 16:36:12 for Cutlass by   +Generated on Wed Sep 19 2018 13:58:50 for Cutlass by   doxygen 1.8.14
    diff --git a/docs/hgemm__traits_8h.html b/docs/hgemm__traits_8h.html index 283ceb75..bb8e72d9 100644 --- a/docs/hgemm__traits_8h.html +++ b/docs/hgemm__traits_8h.html @@ -82,23 +82,23 @@ $(function() {

    Defies structural properties of half-precision GEMM computation. More...

    -
    #include <cutlass/convert.h>
    -#include <cutlass/reshape_tile.h>
    -#include <cutlass/gemm/gemm.h>
    -#include <cutlass/gemm/gemm_epilogue.h>
    -#include <cutlass/gemm/gemm_epilogue_traits.h>
    -#include <cutlass/gemm/gemm_global_tile.h>
    -#include <cutlass/gemm/gemm_shared_tile.h>
    -#include <cutlass/gemm/gemm_traits.h>
    -#include <cutlass/gemm/hgemm_global_tile.h>
    -#include <cutlass/gemm/hgemm_multiply_add.h>
    -#include <cutlass/gemm/hgemm_swizzle.h>
    +
    #include "cutlass/convert.h"
    +#include "cutlass/reshape_tile.h"
    +#include "cutlass/gemm/gemm.h"
    +#include "cutlass/gemm/gemm_epilogue.h"
    +#include "cutlass/gemm/gemm_epilogue_traits.h"
    +#include "cutlass/gemm/gemm_global_tile.h"
    +#include "cutlass/gemm/gemm_shared_tile.h"
    +#include "cutlass/gemm/gemm_traits.h"
    +#include "cutlass/gemm/hgemm_global_tile.h"
    +#include "cutlass/gemm/hgemm_multiply_add.h"
    +#include "cutlass/gemm/hgemm_swizzle.h"

    Go to the source code of this file.

    - + @@ -120,9 +120,9 @@ Classes - + - +

    Classes

    struct  cutlass::gemm::HgemmConfig< OutputTile_, AccumulatorsPerThread_, kScalarsPerLdgA_, kScalarsPerLdgB_ >
    struct  cutlass::gemm::HgemmConfig< OutputTile_, ThreadGemmShape_, kScalarsPerLdgA_, kScalarsPerLdgB_ >
     
    struct  cutlass::gemm::HgemmTransformerA< kLayout_, Iterator_ >
     
     
    struct  cutlass::gemm::HgemmTileTraitsHelperB< MatrixLayout::kColumnMajor, GemmConfig_ >
     
    struct  cutlass::gemm::HgemmTraitsHelper< kLayoutA_, kLayoutB_, OutputTile_, EpilogueFunctor_, AccumulatorsPerThread_, kScalarsPerLdgA_, kScalarsPerLdgB_, Index_ >
    struct  cutlass::gemm::HgemmTraitsHelper< kLayoutA_, kLayoutB_, OutputTile_, EpilogueFunctor_, ThreadGemmShape_, kScalarsPerLdgA_, kScalarsPerLdgB_, Index_ >
     
    struct  cutlass::gemm::HgemmTraits< kLayoutA_, kLayoutB_, OutputTile_, EpilogueFunctor_, AccumulatorsPerThread_, kScalarsPerLdgA_, kScalarsPerLdgB_, Index_, Helper_ >
    struct  cutlass::gemm::HgemmTraits< kLayoutA_, kLayoutB_, OutputTile_, EpilogueFunctor_, ThreadGemmShape_, kScalarsPerLdgA_, kScalarsPerLdgB_, Index_, Helper_ >
     
    -Generated on Thu May 3 2018 16:36:12 for Cutlass by   +Generated on Wed Sep 19 2018 13:58:50 for Cutlass by   doxygen 1.8.14
    diff --git a/docs/hgemm__traits_8h_source.html b/docs/hgemm__traits_8h_source.html index 0d12493e..db1554c8 100644 --- a/docs/hgemm__traits_8h_source.html +++ b/docs/hgemm__traits_8h_source.html @@ -76,89 +76,87 @@ $(function() {
    hgemm_traits.h
    -Go to the documentation of this file.
    1 /***************************************************************************************************
    2  * Copyright (c) 2017-2018, NVIDIA CORPORATION. All rights reserved.
    3  *
    4  * Redistribution and use in source and binary forms, with or without modification, are permitted
    5  * provided that the following conditions are met:
    6  * * Redistributions of source code must retain the above copyright notice, this list of
    7  * conditions and the following disclaimer.
    8  * * Redistributions in binary form must reproduce the above copyright notice, this list of
    9  * conditions and the following disclaimer in the documentation and/or other materials
    10  * provided with the distribution.
    11  * * Neither the name of the NVIDIA CORPORATION nor the names of its contributors may be used
    12  * to endorse or promote products derived from this software without specific prior written
    13  * permission.
    14  *
    15  * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND ANY EXPRESS OR
    16  * IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND
    17  * FITNESS FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE
    18  * FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING,
    19  * BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS;
    20  * OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT,
    21  * STRICT LIABILITY, OR TOR (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
    22  * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
    23  *
    24  **************************************************************************************************/
    28 #pragma once
    29 
    30 #include <cutlass/convert.h>
    31 #include <cutlass/reshape_tile.h>
    32 
    33 #include <cutlass/gemm/gemm.h>
    34 #include <cutlass/gemm/gemm_epilogue.h>
    35 #include <cutlass/gemm/gemm_epilogue_traits.h>
    36 #include <cutlass/gemm/gemm_global_tile.h>
    37 #include <cutlass/gemm/gemm_shared_tile.h>
    38 #include <cutlass/gemm/gemm_traits.h>
    39 #include <cutlass/gemm/hgemm_global_tile.h>
    40 #include <cutlass/gemm/hgemm_multiply_add.h>
    41 #include <cutlass/gemm/hgemm_swizzle.h>
    42 
    43 namespace cutlass {
    44 namespace gemm {
    45 
    47 
    48 template <
    50  typename OutputTile_,
    52  typename AccumulatorsPerThread_,
    54  int kScalarsPerLdgA_ = 2,
    56  int kScalarsPerLdgB_ = 2>
    57 struct HgemmConfig
    58  : public GemmConfig<
    60  half,
    62  half,
    64  half,
    66  half,
    68  OutputTile_,
    70  ThreadMultiplyAdd<AccumulatorsPerThread_, Shape<1, 4, 8>, half, half, half>,
    72  kScalarsPerLdgA_,
    74  kScalarsPerLdgA_,
    76  8,
    78  kScalarsPerLdgB_,
    80  kScalarsPerLdgB_,
    82  8,
    84  2,
    86  8,
    88  2,
    90  2> {};
    91 
    93 
    94 template <enum MatrixLayout::Kind kLayout_, typename Iterator_>
    95 struct HgemmTransformerA {};
    96 
    97 template <typename Iterator_>
    98 struct HgemmTransformerA<MatrixLayout::kColumnMajor, Iterator_> {
    99  typedef Convert<typename Iterator_::Fragment, typename Iterator_::Fragment> Transformer;
    100 };
    101 
    102 template <typename Iterator_>
    103 struct HgemmTransformerA<MatrixLayout::kRowMajor, Iterator_> {
    104  typedef HgemmSwizzle<Iterator_> Transformer;
    105 };
    106 
    108 
    109 template <enum MatrixLayout::Kind kLayout_, typename Iterator_>
    110 struct HgemmTransformerB {};
    111 
    112 template <typename Iterator_>
    113 struct HgemmTransformerB<MatrixLayout::kRowMajor, Iterator_> {
    114  typedef Convert<typename Iterator_::Fragment, typename Iterator_::Fragment> Transformer;
    115 };
    116 
    117 template <typename Iterator_>
    118 struct HgemmTransformerB<MatrixLayout::kColumnMajor, Iterator_> {
    119  typedef HgemmSwizzle<Iterator_> Transformer;
    120 };
    121 
    123 
    124 template <enum MatrixLayout::Kind kLayout_, typename GemmConfig_>
    125 struct HgemmTileTraitsHelperA : public GemmTileTraitsHelperA<kLayout_, GemmConfig_> {};
    126 
    128 
    129 template <typename GemmConfig_>
    130 struct HgemmTileTraitsHelperA<MatrixLayout::kRowMajor, GemmConfig_>
    131  : public GemmTileTraitsHelperA<MatrixLayout::kRowMajor, GemmConfig_> {
    133  typedef GemmTileTraitsHelperA<MatrixLayout::kRowMajor, GemmConfig_> Base;
    134 
    136  typedef HgemmCrosswiseGlobalTileTraits<
    137  GemmOperand::kA,
    138  // The layout.
    139  MatrixLayout::kRowMajor,
    140  // The pointer.
    141  half const,
    142  // The tile has size MxK in GEMM's terminology.
    143  Shape<1, GemmConfig_::OutputTile::kW, GemmConfig_::OutputTile::kD>,
    144  // The threads are distributed as (threads / K ) x K (the traits may reorganize).
    145  Shape<1, GemmConfig_::kThreads / GemmConfig_::OutputTile::kD, GemmConfig_::OutputTile::kD>,
    146  // The number of scalars per LDG (LDG.32 or LDG.128, etc)
    147  GemmConfig_::kScalarsPerLdgA>
    148  GlobalTileTraits;
    149 
    151  typedef GemmSharedStoreWithSkewTileAbTraits<
    152  // The pointer.
    153  half,
    154  // The tile has size KxM in GEMM's terminology.
    155  Shape<GemmConfig_::kStages,
    156  GemmConfig_::OutputTile::kD / GemmConfig_::InstructionShape::kD,
    157  GemmConfig_::OutputTile::kW * GemmConfig_::InstructionShape::kD>,
    158  // The threads are distributed as warps x 32(the traits may reorganize).
    159  typename GlobalTileTraits::Threads,
    160  // The number of scalars per STS (STS.32 or STS.128, etc).
    161  2,
    162  // The skew to avoid bank conflicts added in the tile W dimension.
    163  128 / sizeof(half) / GlobalTileTraits::Threads::kW / 2>
    164  SharedStoreTileTraits;
    165 
    167  typedef GemmSharedLoadTileATraits<
    168  // The pointer.
    169  half const,
    170  // The output tile size.
    171  typename GemmConfig_::OutputTile,
    172  // The number of warps.
    173  typename GemmConfig_::Warps,
    174  // The number of threads per warp.
    175  typename GemmConfig_::MultiplyAdd::ThreadsPerWarp,
    176  // The shape of the FMA instruction.
    177  typename GemmConfig_::InstructionShape,
    178  // The number of stages.
    179  GemmConfig_::kStages,
    180  // The number of scalars per LDS.
    181  8,
    182  // The skew.
    183  SharedStoreTileTraits::kSkew>
    184  SharedLoadTileTraits;
    185 };
    186 
    188 
    189 template <enum MatrixLayout::Kind kLayout_, typename GemmConfig_>
    190 struct HgemmTileTraitsHelperB : public GemmTileTraitsHelperB<kLayout_, GemmConfig_> {};
    191 
    193 
    194 template <typename GemmConfig_>
    195 struct HgemmTileTraitsHelperB<MatrixLayout::kColumnMajor, GemmConfig_>
    196  : public GemmTileTraitsHelperB<MatrixLayout::kColumnMajor, GemmConfig_> {
    198  typedef GemmTileTraitsHelperB<MatrixLayout::kColumnMajor, GemmConfig_> Base;
    199 
    201  typedef HgemmCrosswiseGlobalTileTraits<
    202  GemmOperand::kB,
    203  // The layout.
    204  MatrixLayout::kColumnMajor,
    205  // The pointer.
    206  half const,
    207  // The tile has size KxN in GEMM's terminology.
    208  Shape<1, GemmConfig_::OutputTile::kH, GemmConfig_::OutputTile::kD>,
    209  // The threads are distributed as (threads / K) x K (the traits may reorganize).
    210  Shape<1, GemmConfig_::kThreads / GemmConfig_::OutputTile::kD, GemmConfig_::OutputTile::kD>,
    211  // The number of scalars per LDG (LDG.32 or LDG.128, etc)
    212  GemmConfig_::kScalarsPerLdgB>
    213  GlobalTileTraits;
    214 
    216  typedef GemmSharedStoreWithSkewTileAbTraits<
    217  // The pointer.
    218  half,
    219  // The tile has size KxN in GEMM's terminology.
    220  Shape<GemmConfig_::kStages,
    221  GemmConfig_::OutputTile::kD / GemmConfig_::InstructionShape::kD,
    222  GemmConfig_::OutputTile::kH * GemmConfig_::InstructionShape::kD>,
    223  // The threads are distributed as (threads / K) x K (the traits may reorganize).
    224  typename GlobalTileTraits::Threads,
    225  // The number of scalars per STS (STS.32 or STS.128, etc).
    226  2,
    227  // The skew to avoid bank conflicts added in the tile W dimension.
    228  128 / sizeof(half) / GlobalTileTraits::Threads::kW / 2>
    229  SharedStoreTileTraits;
    230 
    232  typedef GemmSharedLoadTileBTraits<
    233  // The pointer.
    234  half const,
    235  // The output tile size.
    236  typename GemmConfig_::OutputTile,
    237  // The number of warps.
    238  typename GemmConfig_::Warps,
    239  // The number of threads per warp.
    240  typename GemmConfig_::MultiplyAdd::ThreadsPerWarp,
    241  // The shape of the FMA instruction.
    242  typename GemmConfig_::InstructionShape,
    243  // The number of stages.
    244  GemmConfig_::kStages,
    245  // The number of scalars per LDS.
    246  8,
    247  // The skew.
    248  SharedStoreTileTraits::kSkew>
    249  SharedLoadTileTraits;
    250 };
    251 
    253 
    254 template <
    256  MatrixLayout::Kind kLayoutA_,
    258  MatrixLayout::Kind kLayoutB_,
    260  typename OutputTile_,
    262  typename EpilogueFunctor_,
    264  typename AccumulatorsPerThread_ = Shape<32, 8, 8>,
    266  int kScalarsPerLdgA_ = 2,
    268  int kScalarsPerLdgB_ = 2,
    270  typename Index_ = int>
    271 struct HgemmTraitsHelper {
    273  typedef HgemmConfig<OutputTile_, AccumulatorsPerThread_, kScalarsPerLdgA_, kScalarsPerLdgB_>
    274  GemmConfig;
    276  typedef HgemmTileTraitsHelperA<kLayoutA_, GemmConfig> GemmTileTraitsHelperA;
    278  typedef HgemmTileTraitsHelperB<kLayoutB_, GemmConfig> GemmTileTraitsHelperB;
    279 
    281  typedef GemmGlobalIteratorAb<typename GemmTileTraitsHelperA::GlobalTileTraits, Index_>
    282  GlobalLoadIteratorA;
    284  typedef typename HgemmTransformerA<GemmTileTraitsHelperA::kLayout,
    285  GlobalLoadIteratorA>::Transformer GlobalTransformerA;
    287  typedef TileStoreIterator<typename GemmTileTraitsHelperA::SharedStoreTileTraits,
    288  typename GemmTileTraitsHelperA::SharedStoreTileTraits::Scalar,
    289  IteratorAdvance::kH,
    290  MemorySpace::kShared>
    291  SharedStoreIteratorA;
    293  typedef GlobalLoadStream<GlobalLoadIteratorA, SharedStoreIteratorA, GlobalTransformerA>
    294  GlobalLoadStreamA;
    295 
    297  typedef GemmGlobalIteratorAb<typename GemmTileTraitsHelperB::GlobalTileTraits, Index_>
    298  GlobalLoadIteratorB;
    299  // The default transformer for B.
    300  typedef typename HgemmTransformerB<GemmTileTraitsHelperB::kLayout,
    301  GlobalLoadIteratorB>::Transformer GlobalTransformerB;
    303  typedef TileStoreIterator<typename GemmTileTraitsHelperB::SharedStoreTileTraits,
    304  typename GemmTileTraitsHelperB::SharedStoreTileTraits::Scalar,
    305  IteratorAdvance::kH,
    306  MemorySpace::kShared>
    307  SharedStoreIteratorB;
    309  typedef GlobalLoadStream<GlobalLoadIteratorB, SharedStoreIteratorB, GlobalTransformerB>
    310  GlobalLoadStreamB;
    311 
    313  typedef TileLoadIterator<typename GemmTileTraitsHelperA::SharedLoadTileTraits,
    314  typename GemmTileTraitsHelperA::SharedLoadTileTraits::Scalar,
    315  IteratorAdvance::kH,
    316  MemorySpace::kShared>
    317  SharedLoadIteratorA;
    319  typedef SharedLoadStream<SharedLoadIteratorA> SharedLoadStreamA;
    321  typedef TileLoadIterator<typename GemmTileTraitsHelperB::SharedLoadTileTraits,
    322  typename GemmTileTraitsHelperB::SharedLoadTileTraits::Scalar,
    323  IteratorAdvance::kH,
    324  MemorySpace::kShared>
    325  SharedLoadIteratorB;
    327  typedef SharedLoadStream<SharedLoadIteratorB> SharedLoadStreamB;
    328 
    330  typedef typename GemmConfig::MultiplyAdd MultiplyAdd;
    332  typedef ClearAccumulators<typename MultiplyAdd::ScalarC> ClearAccumulators;
    333 
    335  typedef SimplifiedGemmEpilogueTraits<GemmConfig, EpilogueFunctor_, Index_> GemmEpilogueTraits;
    337  typedef GemmEpilogue<GemmEpilogueTraits> Epilogue;
    338 };
    339 
    341 
    342 template <
    344  MatrixLayout::Kind kLayoutA_,
    346  MatrixLayout::Kind kLayoutB_,
    348  typename OutputTile_ = Shape<8, 128, 128>,
    350  typename EpilogueFunctor_ = LinearScaling<half>,
    352  typename AccumulatorsPerThread_ = Shape<8, 8, 16>,
    354  int kScalarsPerLdgA_ = 2,
    356  int kScalarsPerLdgB_ = 2,
    358  typename Index_ = int,
    360  typename Helper_ = HgemmTraitsHelper<kLayoutA_,
    361  kLayoutB_,
    362  OutputTile_,
    363  EpilogueFunctor_,
    364  AccumulatorsPerThread_,
    365  kScalarsPerLdgA_,
    366  kScalarsPerLdgB_,
    367  Index_> >
    368 struct HgemmTraits : public GemmTraits<
    369  // The config.
    370  typename Helper_::GemmConfig,
    371  // The stream to load A from global memory to shared memory.
    372  typename Helper_::GlobalLoadStreamA,
    373  // The stream to load B from global memory to shared memory.
    374  typename Helper_::GlobalLoadStreamB,
    375  // The stream to load A from shared memory.
    376  typename Helper_::SharedLoadStreamA,
    377  // The stream to load B from shared memory.
    378  typename Helper_::SharedLoadStreamB,
    379  // The epilogue.
    380  typename Helper_::Epilogue,
    381  // The block swizzle to reorganize the grid.
    382  IdentityBlockSwizzle,
    383  // The index.
    384  Index_,
    385  // The tool used to clear accumulators.
    386  typename Helper_::ClearAccumulators> {};
    387 
    389 
    390 } // namespace gemm
    391 } // namespace cutlass
    cutlass::gemm::HgemmTraitsHelper::GlobalLoadIteratorA
    GemmGlobalIteratorAb< typename GemmTileTraitsHelperA::GlobalTileTraits, Index_ > GlobalLoadIteratorA
    The iterator to load A from global memory.
    Definition: hgemm_traits.h:282
    -
    cutlass::MemorySpace::kShared
    Definition: load_store.h:42
    -
    cutlass::gemm::HgemmTransformerB< MatrixLayout::kColumnMajor, Iterator_ >::Transformer
    HgemmSwizzle< Iterator_ > Transformer
    Definition: hgemm_traits.h:119
    +Go to the documentation of this file.
    1 /***************************************************************************************************
    2  * Copyright (c) 2017-2018, NVIDIA CORPORATION. All rights reserved.
    3  *
    4  * Redistribution and use in source and binary forms, with or without modification, are permitted
    5  * provided that the following conditions are met:
    6  * * Redistributions of source code must retain the above copyright notice, this list of
    7  * conditions and the following disclaimer.
    8  * * Redistributions in binary form must reproduce the above copyright notice, this list of
    9  * conditions and the following disclaimer in the documentation and/or other materials
    10  * provided with the distribution.
    11  * * Neither the name of the NVIDIA CORPORATION nor the names of its contributors may be used
    12  * to endorse or promote products derived from this software without specific prior written
    13  * permission.
    14  *
    15  * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND ANY EXPRESS OR
    16  * IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND
    17  * FITNESS FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE
    18  * FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING,
    19  * BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS;
    20  * OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT,
    21  * STRICT LIABILITY, OR TOR (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
    22  * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
    23  *
    24  **************************************************************************************************/
    28 #pragma once
    29 
    30 #include "cutlass/convert.h"
    31 #include "cutlass/reshape_tile.h"
    32 
    33 #include "cutlass/gemm/gemm.h"
    34 #include "cutlass/gemm/gemm_epilogue.h"
    35 #include "cutlass/gemm/gemm_epilogue_traits.h"
    36 #include "cutlass/gemm/gemm_global_tile.h"
    37 #include "cutlass/gemm/gemm_shared_tile.h"
    38 #include "cutlass/gemm/gemm_traits.h"
    39 #include "cutlass/gemm/hgemm_global_tile.h"
    40 #include "cutlass/gemm/hgemm_multiply_add.h"
    41 #include "cutlass/gemm/hgemm_swizzle.h"
    42 
    43 namespace cutlass {
    44 namespace gemm {
    45 
    47 
    48 template <
    50  typename OutputTile_,
    52  typename ThreadGemmShape_,
    54  int kScalarsPerLdgA_ = 2,
    56  int kScalarsPerLdgB_ = 2>
    57 struct HgemmConfig : public GemmConfig<
    59  half,
    61  half,
    63  half,
    65  half,
    67  OutputTile_,
    69  ThreadMultiplyAdd<ThreadGemmShape_, Shape<1, 4, 8>, half, half, half>,
    71  kScalarsPerLdgA_,
    73  kScalarsPerLdgA_,
    75  8,
    77  kScalarsPerLdgB_,
    79  kScalarsPerLdgB_,
    81  8,
    83  2,
    85  8,
    87  2,
    89  2,
    91  false,
    93  true,
    95  false
    96  > {};
    97 
    99 
    100 template <enum MatrixLayout::Kind kLayout_, typename Iterator_>
    101 struct HgemmTransformerA {};
    102 
    103 template <typename Iterator_>
    104 struct HgemmTransformerA<MatrixLayout::kColumnMajor, Iterator_> {
    105  typedef Convert<typename Iterator_::Fragment, typename Iterator_::Fragment> Transformer;
    106 };
    107 
    108 template <typename Iterator_>
    109 struct HgemmTransformerA<MatrixLayout::kRowMajor, Iterator_> {
    110  typedef HgemmSwizzle<Iterator_> Transformer;
    111 };
    112 
    114 
    115 template <enum MatrixLayout::Kind kLayout_, typename Iterator_>
    116 struct HgemmTransformerB {};
    117 
    118 template <typename Iterator_>
    119 struct HgemmTransformerB<MatrixLayout::kRowMajor, Iterator_> {
    120  typedef Convert<typename Iterator_::Fragment, typename Iterator_::Fragment> Transformer;
    121 };
    122 
    123 template <typename Iterator_>
    124 struct HgemmTransformerB<MatrixLayout::kColumnMajor, Iterator_> {
    125  typedef HgemmSwizzle<Iterator_> Transformer;
    126 };
    127 
    129 
    130 template <enum MatrixLayout::Kind kLayout_, typename GemmConfig_>
    131 struct HgemmTileTraitsHelperA : public GemmTileTraitsHelperA<kLayout_, GemmConfig_> {};
    132 
    134 
    135 template <typename GemmConfig_>
    136 struct HgemmTileTraitsHelperA<MatrixLayout::kRowMajor, GemmConfig_>
    137  : public GemmTileTraitsHelperA<MatrixLayout::kRowMajor, GemmConfig_> {
    139  typedef GemmTileTraitsHelperA<MatrixLayout::kRowMajor, GemmConfig_> Base;
    140 
    142  typedef HgemmCrosswiseGlobalTileTraits<
    143  GemmOperand::kA,
    144  // The layout.
    145  MatrixLayout::kRowMajor,
    146  // The pointer.
    147  half const,
    148  // The tile has size MxK in GEMM's terminology.
    149  Shape<1, GemmConfig_::OutputTile::kW, GemmConfig_::OutputTile::kD>,
    150  // The threads are distributed as (threads / K ) x K (the traits may reorganize).
    151  Shape<1, GemmConfig_::kThreads / GemmConfig_::OutputTile::kD, GemmConfig_::OutputTile::kD>,
    152  // The number of scalars per LDG (LDG.32 or LDG.128, etc)
    153  GemmConfig_::kScalarsPerLdgA>
    154  GlobalTileTraits;
    155 
    156  static int const kSkewA = 128 / sizeof(half) / GlobalTileTraits::Threads::kW / 2;
    157 
    159  typedef GemmSharedStoreWithSkewTileAbTraits <
    160  // The pointer.
    161  half,
    162  // The tile has size KxM in GEMM's terminology.
    163  Shape<GemmConfig_::kStages,
    164  GemmConfig_::OutputTile::kD / GemmConfig_::InstructionShape::kD,
    165  GemmConfig_::OutputTile::kW * GemmConfig_::InstructionShape::kD>,
    166  // The threads are distributed as warps x 32(the traits may reorganize).
    167  typename GlobalTileTraits::Threads,
    168  // The number of scalars per STS (STS.32 or STS.128, etc).
    169  2,
    170  // The skew to avoid bank conflicts added in the tile W dimension.
    171  kSkewA<GemmConfig_::kScalarsPerLdsA ? GemmConfig_::kScalarsPerLdsA : kSkewA>
    172  SharedStoreTileTraits;
    173 
    175  typedef GemmSharedLoadTileATraits<
    176  // The pointer.
    177  half const,
    178  // The output tile size.
    179  typename GemmConfig_::OutputTile,
    180  // The number of warps.
    181  typename GemmConfig_::Warps,
    182  // The number of threads per warp.
    183  typename GemmConfig_::MultiplyAdd::ThreadsPerWarp,
    184  // The shape of the FMA instruction.
    185  typename GemmConfig_::InstructionShape,
    186  // The number of stages.
    187  GemmConfig_::kStages,
    188  // The number of scalars per LDS.
    189  8,
    190  // The skew.
    191  SharedStoreTileTraits::kSkew>
    192  SharedLoadTileTraits;
    193 };
    194 
    196 
    197 template <enum MatrixLayout::Kind kLayout_, typename GemmConfig_>
    198 struct HgemmTileTraitsHelperB : public GemmTileTraitsHelperB<kLayout_, GemmConfig_> {};
    199 
    201 
    202 template <typename GemmConfig_>
    203 struct HgemmTileTraitsHelperB<MatrixLayout::kColumnMajor, GemmConfig_>
    204  : public GemmTileTraitsHelperB<MatrixLayout::kColumnMajor, GemmConfig_> {
    206  typedef GemmTileTraitsHelperB<MatrixLayout::kColumnMajor, GemmConfig_> Base;
    207 
    209  typedef HgemmCrosswiseGlobalTileTraits<
    210  GemmOperand::kB,
    211  // The layout.
    212  MatrixLayout::kColumnMajor,
    213  // The pointer.
    214  half const,
    215  // The tile has size KxN in GEMM's terminology.
    216  Shape<1, GemmConfig_::OutputTile::kH, GemmConfig_::OutputTile::kD>,
    217  // The threads are distributed as (threads / K) x K (the traits may reorganize).
    218  Shape<1, GemmConfig_::kThreads / GemmConfig_::OutputTile::kD, GemmConfig_::OutputTile::kD>,
    219  // The number of scalars per LDG (LDG.32 or LDG.128, etc)
    220  GemmConfig_::kScalarsPerLdgB>
    221  GlobalTileTraits;
    222 
    223  static int const kSkewB = 128 / sizeof(half) / GlobalTileTraits::Threads::kW / 2;
    224 
    226  typedef GemmSharedStoreWithSkewTileAbTraits <
    227  // The pointer.
    228  half,
    229  // The tile has size KxN in GEMM's terminology.
    230  Shape<GemmConfig_::kStages,
    231  GemmConfig_::OutputTile::kD / GemmConfig_::InstructionShape::kD,
    232  GemmConfig_::OutputTile::kH * GemmConfig_::InstructionShape::kD>,
    233  // The threads are distributed as (threads / K) x K (the traits may reorganize).
    234  typename GlobalTileTraits::Threads,
    235  // The number of scalars per STS (STS.32 or STS.128, etc).
    236  2,
    237  // The skew to avoid bank conflicts added in the tile W dimension.
    238  kSkewB<GemmConfig_::kScalarsPerLdsB ? GemmConfig_::kScalarsPerLdsB : kSkewB>
    239  SharedStoreTileTraits;
    240 
    242  typedef GemmSharedLoadTileBTraits<
    243  // The pointer.
    244  half const,
    245  // The output tile size.
    246  typename GemmConfig_::OutputTile,
    247  // The number of warps.
    248  typename GemmConfig_::Warps,
    249  // The number of threads per warp.
    250  typename GemmConfig_::MultiplyAdd::ThreadsPerWarp,
    251  // The shape of the FMA instruction.
    252  typename GemmConfig_::InstructionShape,
    253  // The number of stages.
    254  GemmConfig_::kStages,
    255  // The number of scalars per LDS.
    256  8,
    257  // The skew.
    258  SharedStoreTileTraits::kSkew>
    259  SharedLoadTileTraits;
    260 };
    261 
    263 
    264 template <
    266  MatrixLayout::Kind kLayoutA_,
    268  MatrixLayout::Kind kLayoutB_,
    270  typename OutputTile_,
    272  typename EpilogueFunctor_,
    274  typename ThreadGemmShape_,
    276  int kScalarsPerLdgA_ = 2,
    278  int kScalarsPerLdgB_ = 2,
    280  typename Index_ = int>
    281 struct HgemmTraitsHelper {
    283  typedef HgemmConfig<OutputTile_, ThreadGemmShape_, kScalarsPerLdgA_, kScalarsPerLdgB_> GemmConfig;
    285  typedef HgemmTileTraitsHelperA<kLayoutA_, GemmConfig> GemmTileTraitsHelperA;
    287  typedef HgemmTileTraitsHelperB<kLayoutB_, GemmConfig> GemmTileTraitsHelperB;
    288 
    290  typedef GemmGlobalIteratorAb<typename GemmTileTraitsHelperA::GlobalTileTraits, Index_>
    291  GlobalLoadIteratorA;
    293  typedef typename HgemmTransformerA<GemmTileTraitsHelperA::kLayout,
    294  GlobalLoadIteratorA>::Transformer GlobalTransformerA;
    296  typedef TileStoreIterator<typename GemmTileTraitsHelperA::SharedStoreTileTraits,
    297  typename GemmTileTraitsHelperA::SharedStoreTileTraits::Scalar,
    298  IteratorAdvance::kH,
    299  MemorySpace::kShared>
    300  SharedStoreIteratorA;
    302  typedef GlobalLoadStream<GemmOperand::kA,
    303  GlobalLoadIteratorA,
    304  SharedStoreIteratorA,
    305  GlobalTransformerA>
    306  GlobalLoadStreamA;
    307 
    309  typedef GemmGlobalIteratorAb<typename GemmTileTraitsHelperB::GlobalTileTraits, Index_>
    310  GlobalLoadIteratorB;
    311  // The default transformer for B.
    312  typedef typename HgemmTransformerB<GemmTileTraitsHelperB::kLayout,
    313  GlobalLoadIteratorB>::Transformer GlobalTransformerB;
    315  typedef TileStoreIterator<typename GemmTileTraitsHelperB::SharedStoreTileTraits,
    316  typename GemmTileTraitsHelperB::SharedStoreTileTraits::Scalar,
    317  IteratorAdvance::kH,
    318  MemorySpace::kShared>
    319  SharedStoreIteratorB;
    321  typedef GlobalLoadStream<GemmOperand::kB,
    322  GlobalLoadIteratorB,
    323  SharedStoreIteratorB,
    324  GlobalTransformerB>
    325  GlobalLoadStreamB;
    326 
    328  typedef TileLoadIterator<typename GemmTileTraitsHelperA::SharedLoadTileTraits,
    329  typename GemmTileTraitsHelperA::SharedLoadTileTraits::Scalar,
    330  IteratorAdvance::kH,
    331  MemorySpace::kShared>
    332  SharedLoadIteratorA;
    334  typedef SharedLoadStream<SharedLoadIteratorA> SharedLoadStreamA;
    336  typedef TileLoadIterator<typename GemmTileTraitsHelperB::SharedLoadTileTraits,
    337  typename GemmTileTraitsHelperB::SharedLoadTileTraits::Scalar,
    338  IteratorAdvance::kH,
    339  MemorySpace::kShared>
    340  SharedLoadIteratorB;
    342  typedef SharedLoadStream<SharedLoadIteratorB> SharedLoadStreamB;
    343 
    345  typedef typename GemmConfig::MultiplyAdd MultiplyAdd;
    347  typedef ClearAccumulators<typename MultiplyAdd::ScalarC> ClearAccumulators;
    348 
    350  typedef SimplifiedGemmEpilogueTraits<GemmConfig, EpilogueFunctor_, Index_> GemmEpilogueTraits;
    352  typedef GemmEpilogue<GemmEpilogueTraits> Epilogue;
    353 };
    354 
    356 
    357 template <
    359  MatrixLayout::Kind kLayoutA_,
    361  MatrixLayout::Kind kLayoutB_,
    363  typename OutputTile_ = Shape<8, 128, 128>,
    365  typename EpilogueFunctor_ = LinearScaling<half>,
    367  typename ThreadGemmShape_ = Shape<8, 8, 16>,
    369  int kScalarsPerLdgA_ = 2,
    371  int kScalarsPerLdgB_ = 2,
    373  typename Index_ = int,
    375  typename Helper_ = HgemmTraitsHelper<kLayoutA_,
    376  kLayoutB_,
    377  OutputTile_,
    378  EpilogueFunctor_,
    379  ThreadGemmShape_,
    380  kScalarsPerLdgA_,
    381  kScalarsPerLdgB_,
    382  Index_> >
    383 struct HgemmTraits : public GemmTraits<
    384  // The config.
    385  typename Helper_::GemmConfig,
    386  // The stream to load A from global memory to shared memory.
    387  typename Helper_::GlobalLoadStreamA,
    388  // The stream to load B from global memory to shared memory.
    389  typename Helper_::GlobalLoadStreamB,
    390  // The stream to load A from shared memory.
    391  typename Helper_::SharedLoadStreamA,
    392  // The stream to load B from shared memory.
    393  typename Helper_::SharedLoadStreamB,
    394  // The epilogue.
    395  typename Helper_::Epilogue,
    396  // The block swizzle to reorganize the grid.
    397  IdentityBlockSwizzle,
    398  // The index.
    399  Index_,
    400  // The tool used to clear accumulators.
    401  typename Helper_::ClearAccumulators> {};
    402 
    404 
    405 } // namespace gemm
    406 } // namespace cutlass
    cutlass::gemm::HgemmTraitsHelper::SharedLoadStreamB
    SharedLoadStream< SharedLoadIteratorB > SharedLoadStreamB
    The stream to load B from shared memory.
    Definition: hgemm_traits.h:342
    +
    cutlass::gemm::HgemmTraitsHelper::GlobalLoadIteratorB
    GemmGlobalIteratorAb< typename GemmTileTraitsHelperB::GlobalTileTraits, Index_ > GlobalLoadIteratorB
    The iterator to load B from global memory.
    Definition: hgemm_traits.h:310
    +
    cutlass::MemorySpace::kShared
    Definition: load_store.h:41
    +
    cutlass::gemm::HgemmTransformerB< MatrixLayout::kColumnMajor, Iterator_ >::Transformer
    HgemmSwizzle< Iterator_ > Transformer
    Definition: hgemm_traits.h:125
    cutlass
    Definition: convert.h:33
    -
    cutlass::gemm::GemmSharedLoadTileATraits
    Definition: gemm_shared_tile.h:129
    +
    cutlass::gemm::HgemmTraitsHelper::GemmConfig
    HgemmConfig< OutputTile_, ThreadGemmShape_, kScalarsPerLdgA_, kScalarsPerLdgB_ > GemmConfig
    The HGEMM config.
    Definition: hgemm_traits.h:283
    +
    cutlass::gemm::GemmSharedLoadTileATraits
    Definition: gemm_shared_tile.h:128
    -
    cutlass::gemm::GemmEpilogue
    Definition: gemm_epilogue.h:53
    +
    cutlass::gemm::GemmEpilogue
    Definition: gemm_epilogue.h:42
    gemm_global_tile.h
    Defines iterators for efficiently loading and storing to global memory.
    -
    cutlass::gemm::HgemmTraitsHelper::GlobalLoadIteratorB
    GemmGlobalIteratorAb< typename GemmTileTraitsHelperB::GlobalTileTraits, Index_ > GlobalLoadIteratorB
    The iterator to load B from global memory.
    Definition: hgemm_traits.h:298
    -
    cutlass::gemm::HgemmTraitsHelper::ClearAccumulators
    ClearAccumulators< typename MultiplyAdd::ScalarC > ClearAccumulators
    The object to clear accumulators.
    Definition: hgemm_traits.h:332
    +
    cutlass::gemm::HgemmTraitsHelper::GemmEpilogueTraits
    SimplifiedGemmEpilogueTraits< GemmConfig, EpilogueFunctor_, Index_ > GemmEpilogueTraits
    The traits class for the epilogue.
    Definition: hgemm_traits.h:350
    gemm_traits.h
    Defines structural properties of complete GEMM computation.
    -
    cutlass::gemm::HgemmTraitsHelper::SharedStoreIteratorA
    TileStoreIterator< typename GemmTileTraitsHelperA::SharedStoreTileTraits, typename GemmTileTraitsHelperA::SharedStoreTileTraits::Scalar, IteratorAdvance::kH, MemorySpace::kShared > SharedStoreIteratorA
    The iterator to store A to shared memory.
    Definition: hgemm_traits.h:291
    -
    cutlass::gemm::HgemmTraitsHelper::GlobalLoadStreamA
    GlobalLoadStream< GlobalLoadIteratorA, SharedStoreIteratorA, GlobalTransformerA > GlobalLoadStreamA
    The stream to load A from global memory to shared memory.
    Definition: hgemm_traits.h:294
    -
    cutlass::gemm::HgemmTileTraitsHelperB< MatrixLayout::kColumnMajor, GemmConfig_ >::GlobalTileTraits
    HgemmCrosswiseGlobalTileTraits< GemmOperand::kB, MatrixLayout::kColumnMajor, half const, Shape< 1, GemmConfig_::OutputTile::kH, GemmConfig_::OutputTile::kD >, Shape< 1, GemmConfig_::kThreads/GemmConfig_::OutputTile::kD, GemmConfig_::OutputTile::kD >, GemmConfig_::kScalarsPerLdgB > GlobalTileTraits
    The traits class to build the iterator to load data from global memory for B^N.
    Definition: hgemm_traits.h:213
    -
    cutlass::gemm::HgemmTransformerA
    Definition: hgemm_traits.h:95
    -
    cutlass::gemm::HgemmTileTraitsHelperB< MatrixLayout::kColumnMajor, GemmConfig_ >::Base
    GemmTileTraitsHelperB< MatrixLayout::kColumnMajor, GemmConfig_ > Base
    The base config.
    Definition: hgemm_traits.h:198
    -
    cutlass::gemm::HgemmTraitsHelper::SharedLoadStreamA
    SharedLoadStream< SharedLoadIteratorA > SharedLoadStreamA
    The stream to load A from shared memory.
    Definition: hgemm_traits.h:319
    -
    cutlass::gemm::HgemmTransformerA< MatrixLayout::kColumnMajor, Iterator_ >::Transformer
    Convert< typename Iterator_::Fragment, typename Iterator_::Fragment > Transformer
    Definition: hgemm_traits.h:99
    -
    cutlass::gemm::HgemmTraits
    Definition: hgemm_traits.h:368
    -
    cutlass::gemm::HgemmTransformerA< MatrixLayout::kRowMajor, Iterator_ >::Transformer
    HgemmSwizzle< Iterator_ > Transformer
    Definition: hgemm_traits.h:104
    -
    cutlass::IteratorAdvance::kH
    Definition: tile_iterator.h:62
    -
    cutlass::gemm::GemmSharedLoadTileBTraits
    Definition: gemm_shared_tile.h:198
    -
    cutlass::gemm::HgemmTraitsHelper::SharedLoadIteratorB
    TileLoadIterator< typename GemmTileTraitsHelperB::SharedLoadTileTraits, typename GemmTileTraitsHelperB::SharedLoadTileTraits::Scalar, IteratorAdvance::kH, MemorySpace::kShared > SharedLoadIteratorB
    The iterator to load B from shared memory.
    Definition: hgemm_traits.h:325
    -
    cutlass::gemm::GemmGlobalIteratorAb
    Definition: gemm_global_tile.h:159
    -
    cutlass::gemm::HgemmTraitsHelper::Epilogue
    GemmEpilogue< GemmEpilogueTraits > Epilogue
    The epilogue.
    Definition: hgemm_traits.h:337
    -
    cutlass::gemm::HgemmTraitsHelper::GlobalTransformerA
    HgemmTransformerA< GemmTileTraitsHelperA::kLayout, GlobalLoadIteratorA >::Transformer GlobalTransformerA
    The default transformer for A.
    Definition: hgemm_traits.h:285
    +
    cutlass::gemm::HgemmTileTraitsHelperB< MatrixLayout::kColumnMajor, GemmConfig_ >::GlobalTileTraits
    HgemmCrosswiseGlobalTileTraits< GemmOperand::kB, MatrixLayout::kColumnMajor, half const, Shape< 1, GemmConfig_::OutputTile::kH, GemmConfig_::OutputTile::kD >, Shape< 1, GemmConfig_::kThreads/GemmConfig_::OutputTile::kD, GemmConfig_::OutputTile::kD >, GemmConfig_::kScalarsPerLdgB > GlobalTileTraits
    The traits class to build the iterator to load data from global memory for B^N.
    Definition: hgemm_traits.h:221
    +
    cutlass::gemm::HgemmTransformerA
    Definition: hgemm_traits.h:101
    +
    cutlass::gemm::HgemmTileTraitsHelperB< MatrixLayout::kColumnMajor, GemmConfig_ >::Base
    GemmTileTraitsHelperB< MatrixLayout::kColumnMajor, GemmConfig_ > Base
    The base config.
    Definition: hgemm_traits.h:206
    +
    cutlass::gemm::HgemmTraitsHelper::Epilogue
    GemmEpilogue< GemmEpilogueTraits > Epilogue
    The epilogue.
    Definition: hgemm_traits.h:352
    +
    cutlass::gemm::HgemmTransformerA< MatrixLayout::kColumnMajor, Iterator_ >::Transformer
    Convert< typename Iterator_::Fragment, typename Iterator_::Fragment > Transformer
    Definition: hgemm_traits.h:105
    +
    cutlass::gemm::HgemmTraitsHelper::GlobalLoadStreamA
    GlobalLoadStream< GemmOperand::kA, GlobalLoadIteratorA, SharedStoreIteratorA, GlobalTransformerA > GlobalLoadStreamA
    The stream to load A from global memory to shared memory.
    Definition: hgemm_traits.h:306
    +
    cutlass::gemm::HgemmTraits
    Definition: hgemm_traits.h:383
    +
    cutlass::gemm::HgemmTransformerA< MatrixLayout::kRowMajor, Iterator_ >::Transformer
    HgemmSwizzle< Iterator_ > Transformer
    Definition: hgemm_traits.h:110
    +
    cutlass::gemm::HgemmTraitsHelper::SharedLoadIteratorB
    TileLoadIterator< typename GemmTileTraitsHelperB::SharedLoadTileTraits, typename GemmTileTraitsHelperB::SharedLoadTileTraits::Scalar, IteratorAdvance::kH, MemorySpace::kShared > SharedLoadIteratorB
    The iterator to load B from shared memory.
    Definition: hgemm_traits.h:340
    +
    cutlass::gemm::HgemmTraitsHelper::GlobalTransformerA
    HgemmTransformerA< GemmTileTraitsHelperA::kLayout, GlobalLoadIteratorA >::Transformer GlobalTransformerA
    The default transformer for A.
    Definition: hgemm_traits.h:294
    +
    cutlass::IteratorAdvance::kH
    Definition: tile_iterator.h:65
    +
    cutlass::gemm::GemmSharedLoadTileBTraits
    Definition: gemm_shared_tile.h:200
    +
    cutlass::gemm::GemmGlobalIteratorAb
    Definition: gemm_global_tile.h:163
    gemm_epilogue.h
    Implements the epilogue phase of the GEMM kernel that efficiently updates global memory with the comp...
    -
    cutlass::gemm::GlobalLoadStream
    Definition: gemm_global_stream.h:161
    -
    cutlass::gemm::GemmTileTraitsHelperB
    Definition: gemm_traits.h:273
    -
    cutlass::gemm::HgemmTileTraitsHelperA
    Definition: hgemm_traits.h:125
    -
    cutlass::MatrixLayout
    Describes layouts of matrices.
    Definition: matrix_traits.h:35
    -
    cutlass::gemm::HgemmTraitsHelper::SharedLoadStreamB
    SharedLoadStream< SharedLoadIteratorB > SharedLoadStreamB
    The stream to load B from shared memory.
    Definition: hgemm_traits.h:327
    -
    cutlass::gemm::HgemmTransformerB
    Definition: hgemm_traits.h:110
    -
    cutlass::gemm::HgemmTileTraitsHelperA< MatrixLayout::kRowMajor, GemmConfig_ >::Base
    GemmTileTraitsHelperA< MatrixLayout::kRowMajor, GemmConfig_ > Base
    The base config.
    Definition: hgemm_traits.h:133
    -
    cutlass::gemm::HgemmTraitsHelper::SharedLoadIteratorA
    TileLoadIterator< typename GemmTileTraitsHelperA::SharedLoadTileTraits, typename GemmTileTraitsHelperA::SharedLoadTileTraits::Scalar, IteratorAdvance::kH, MemorySpace::kShared > SharedLoadIteratorA
    The iterator to load A from shared memory.
    Definition: hgemm_traits.h:317
    -
    cutlass::TileLoadIterator
    An iterator implementing Tile Load Iterator Concept for loading a tile from memory.
    Definition: tile_iterator.h:302
    -
    cutlass::gemm::HgemmTraitsHelper::GemmEpilogueTraits
    SimplifiedGemmEpilogueTraits< GemmConfig, EpilogueFunctor_, Index_ > GemmEpilogueTraits
    The traits class for the epilogue.
    Definition: hgemm_traits.h:335
    +
    cutlass::gemm::GlobalLoadStream
    Definition: gemm_global_stream.h:52
    +
    cutlass::gemm::GemmTileTraitsHelperB
    Definition: gemm_traits.h:191
    +
    cutlass::gemm::HgemmTileTraitsHelperA
    Definition: hgemm_traits.h:131
    +
    cutlass::gemm::HgemmTraitsHelper::GemmTileTraitsHelperA
    HgemmTileTraitsHelperA< kLayoutA_, GemmConfig > GemmTileTraitsHelperA
    The GEMM config for A.
    Definition: hgemm_traits.h:285
    +
    cutlass::MatrixLayout
    Defines data layouts of various matrix formats usable by TensorRef and other classes.
    Definition: matrix_traits.h:156
    +
    cutlass::gemm::HgemmTransformerB
    Definition: hgemm_traits.h:116
    +
    cutlass::gemm::HgemmTileTraitsHelperA< MatrixLayout::kRowMajor, GemmConfig_ >::Base
    GemmTileTraitsHelperA< MatrixLayout::kRowMajor, GemmConfig_ > Base
    The base config.
    Definition: hgemm_traits.h:139
    +
    cutlass::TileLoadIterator
    An iterator implementing Tile Load Iterator Concept for loading a tile from memory.
    Definition: tile_iterator.h:399
    gemm_shared_tile.h
    Defines iterators for efficiently loading and storing tiles to and from shared memory.
    -
    cutlass::MatrixLayout::kRowMajor
    Definition: matrix_traits.h:36
    -
    cutlass::gemm::GemmTileTraitsHelperA< MatrixLayout::kRowMajor, GemmConfig_ >
    Definition: gemm_traits.h:205
    -
    cutlass::gemm::SharedLoadStream
    Definition: gemm_shared_stream.h:44
    +
    cutlass::MatrixLayout::kRowMajor
    Definition: matrix_traits.h:159
    +
    cutlass::gemm::GemmTileTraitsHelperA< MatrixLayout::kRowMajor, GemmConfig_ >
    Definition: gemm_traits.h:120
    +
    cutlass::gemm::SharedLoadStream
    Definition: gemm_shared_stream.h:45
    +
    cutlass::gemm::HgemmTraitsHelper::GlobalTransformerB
    HgemmTransformerB< GemmTileTraitsHelperB::kLayout, GlobalLoadIteratorB >::Transformer GlobalTransformerB
    Definition: hgemm_traits.h:313
    reshape_tile.h
    Defines a type for restructuring a tile.
    +
    cutlass::gemm::HgemmTraitsHelper::ClearAccumulators
    ClearAccumulators< typename MultiplyAdd::ScalarC > ClearAccumulators
    The object to clear accumulators.
    Definition: hgemm_traits.h:347
    hgemm_multiply_add.h
    Specialization implementing multiply-add operation on half-precision floating point fragments...
    -
    cutlass::gemm::GemmConfig
    Definition: gemm_traits.h:79
    +
    cutlass::gemm::GemmConfig
    Definition: gemm_config.h:76
    +
    cutlass::gemm::HgemmTraitsHelper::SharedLoadIteratorA
    TileLoadIterator< typename GemmTileTraitsHelperA::SharedLoadTileTraits, typename GemmTileTraitsHelperA::SharedLoadTileTraits::Scalar, IteratorAdvance::kH, MemorySpace::kShared > SharedLoadIteratorA
    The iterator to load A from shared memory.
    Definition: hgemm_traits.h:332
    hgemm_swizzle.h
    Transposes a tile of 16b elements. Used by HGEMM to construct a K-strided layout in shared memory for...
    -
    cutlass::gemm::GemmTileTraitsHelperA
    Definition: gemm_traits.h:137
    -
    cutlass::gemm::HgemmTileTraitsHelperB< MatrixLayout::kColumnMajor, GemmConfig_ >::SharedLoadTileTraits
    GemmSharedLoadTileBTraits< half const, typename GemmConfig_::OutputTile, typename GemmConfig_::Warps, typename GemmConfig_::MultiplyAdd::ThreadsPerWarp, typename GemmConfig_::InstructionShape, GemmConfig_::kStages, 8, SharedStoreTileTraits::kSkew > SharedLoadTileTraits
    The traits class to build the iterator to load from shared memory for B^N.
    Definition: hgemm_traits.h:249
    -
    cutlass::GemmOperand::kB
    Definition: matrix_traits.h:43
    -
    cutlass::gemm::HgemmTraitsHelper::GemmConfig
    HgemmConfig< OutputTile_, AccumulatorsPerThread_, kScalarsPerLdgA_, kScalarsPerLdgB_ > GemmConfig
    The HGEMM config.
    Definition: hgemm_traits.h:274
    -
    cutlass::gemm::HgemmTileTraitsHelperB
    Definition: hgemm_traits.h:190
    -
    cutlass::gemm::HgemmTraitsHelper::GlobalLoadStreamB
    GlobalLoadStream< GlobalLoadIteratorB, SharedStoreIteratorB, GlobalTransformerB > GlobalLoadStreamB
    The stream to load B from global memory to shared memory.
    Definition: hgemm_traits.h:310
    -
    cutlass::gemm::HgemmTraitsHelper::MultiplyAdd
    GemmConfig::MultiplyAdd MultiplyAdd
    The functor to do the multiply-add in the main loop.
    Definition: hgemm_traits.h:330
    -
    cutlass::gemm::HgemmTraitsHelper::GemmTileTraitsHelperB
    HgemmTileTraitsHelperB< kLayoutB_, GemmConfig > GemmTileTraitsHelperB
    The GEMM config for B.
    Definition: hgemm_traits.h:278
    -
    cutlass::gemm::GemmTraits
    Definition: gemm_traits.h:428
    +
    cutlass::gemm::GemmTileTraitsHelperA
    Definition: gemm_traits.h:52
    +
    cutlass::GemmOperand::kB
    Definition: matrix_traits.h:357
    +
    cutlass::gemm::HgemmTileTraitsHelperB
    Definition: hgemm_traits.h:198
    +
    cutlass::gemm::HgemmTraitsHelper::MultiplyAdd
    GemmConfig::MultiplyAdd MultiplyAdd
    The functor to do the multiply-add in the main loop.
    Definition: hgemm_traits.h:345
    +
    cutlass::gemm::GemmTraits
    Definition: gemm_traits.h:349
    +
    cutlass::gemm::HgemmTraitsHelper::GemmTileTraitsHelperB
    HgemmTileTraitsHelperB< kLayoutB_, GemmConfig > GemmTileTraitsHelperB
    The GEMM config for B.
    Definition: hgemm_traits.h:287
    cutlass::gemm::HgemmCrosswiseGlobalTileTraits
    Definition: hgemm_global_tile.h:48
    cutlass::Shape
    A Shape implementing Layout Concept describing the dimensions of a cube.
    Definition: shape.h:64
    -
    cutlass::gemm::SimplifiedGemmEpilogueTraits
    Definition: gemm_epilogue_traits.h:300
    -
    cutlass::gemm::HgemmTileTraitsHelperA< MatrixLayout::kRowMajor, GemmConfig_ >::SharedLoadTileTraits
    GemmSharedLoadTileATraits< half const, typename GemmConfig_::OutputTile, typename GemmConfig_::Warps, typename GemmConfig_::MultiplyAdd::ThreadsPerWarp, typename GemmConfig_::InstructionShape, GemmConfig_::kStages, 8, SharedStoreTileTraits::kSkew > SharedLoadTileTraits
    The traits class to build the iterator to load from shared memory for A^T.
    Definition: hgemm_traits.h:184
    -
    cutlass::gemm::HgemmTraitsHelper::GemmTileTraitsHelperA
    HgemmTileTraitsHelperA< kLayoutA_, GemmConfig > GemmTileTraitsHelperA
    The GEMM config for A.
    Definition: hgemm_traits.h:276
    -
    cutlass::gemm::ThreadMultiplyAdd
    Template performing matrix multiply-add operation within a thread.
    Definition: thread_multiply_add.h:43
    -
    cutlass::MatrixLayout::kColumnMajor
    Definition: matrix_traits.h:36
    -
    cutlass::MatrixLayout::Kind
    Kind
    Definition: matrix_traits.h:36
    -
    cutlass::gemm::HgemmTraitsHelper::GlobalTransformerB
    HgemmTransformerB< GemmTileTraitsHelperB::kLayout, GlobalLoadIteratorB >::Transformer GlobalTransformerB
    Definition: hgemm_traits.h:301
    -
    cutlass::gemm::GemmTileTraitsHelperB< MatrixLayout::kColumnMajor, GemmConfig_ >
    Definition: gemm_traits.h:278
    -
    cutlass::gemm::HgemmTraitsHelper
    Definition: hgemm_traits.h:271
    -
    cutlass::gemm::HgemmTileTraitsHelperA< MatrixLayout::kRowMajor, GemmConfig_ >::GlobalTileTraits
    HgemmCrosswiseGlobalTileTraits< GemmOperand::kA, MatrixLayout::kRowMajor, half const, Shape< 1, GemmConfig_::OutputTile::kW, GemmConfig_::OutputTile::kD >, Shape< 1, GemmConfig_::kThreads/GemmConfig_::OutputTile::kD, GemmConfig_::OutputTile::kD >, GemmConfig_::kScalarsPerLdgA > GlobalTileTraits
    The traits class to build the iterator to load data from global memory for A^T.
    Definition: hgemm_traits.h:148
    +
    cutlass::gemm::SimplifiedGemmEpilogueTraits
    Definition: gemm_epilogue_traits.h:323
    +
    cutlass::gemm::GemmGlobalTileTraits::Threads
    ReshapeThreads< VectorizedTile, Threads_ >::Threads Threads
    The threads shape.
    Definition: gemm_global_tile.h:88
    +
    cutlass::gemm::ThreadMultiplyAdd
    Template performing matrix multiply-add operation within a thread.
    Definition: thread_multiply_add.h:44
    +
    cutlass::MatrixLayout::kColumnMajor
    Definition: matrix_traits.h:159
    +
    cutlass::MatrixLayout::Kind
    Kind
    Enumeration defining fundamental contiguous layouts.
    Definition: matrix_traits.h:159
    +
    cutlass::gemm::HgemmTraitsHelper::GlobalLoadIteratorA
    GemmGlobalIteratorAb< typename GemmTileTraitsHelperA::GlobalTileTraits, Index_ > GlobalLoadIteratorA
    The iterator to load A from global memory.
    Definition: hgemm_traits.h:291
    +
    cutlass::gemm::GemmTileTraitsHelperB< MatrixLayout::kColumnMajor, GemmConfig_ >
    Definition: gemm_traits.h:196
    +
    cutlass::gemm::HgemmTraitsHelper
    Definition: hgemm_traits.h:281
    +
    cutlass::gemm::HgemmTileTraitsHelperA< MatrixLayout::kRowMajor, GemmConfig_ >::GlobalTileTraits
    HgemmCrosswiseGlobalTileTraits< GemmOperand::kA, MatrixLayout::kRowMajor, half const, Shape< 1, GemmConfig_::OutputTile::kW, GemmConfig_::OutputTile::kD >, Shape< 1, GemmConfig_::kThreads/GemmConfig_::OutputTile::kD, GemmConfig_::OutputTile::kD >, GemmConfig_::kScalarsPerLdgA > GlobalTileTraits
    The traits class to build the iterator to load data from global memory for A^T.
    Definition: hgemm_traits.h:154
    hgemm_global_tile.h
    Tile traits used to construct global tile iterator for HGEMM. This is intended to partition the threa...
    -
    cutlass::gemm::LinearScaling
    Functor to compute linear combination of fragments.
    Definition: linear_scaling.h:40
    +
    cutlass::gemm::LinearScaling
    Functor to compute linear combination of fragments.
    Definition: linear_scaling.h:51
    cutlass::Convert
    Definition: convert.h:38
    -
    cutlass::GemmOperand::kA
    Definition: matrix_traits.h:43
    +
    cutlass::GemmOperand::kA
    Definition: matrix_traits.h:357
    gemm.h
    Implements a software-pipelined efficient GEMM.
    -
    cutlass::gemm::GemmGlobalTileTraits::Threads
    ReshapeThreads< Tile, Threads_ >::Threads Threads
    The threads shape.
    Definition: gemm_global_tile.h:87
    +
    cutlass::gemm::HgemmTraitsHelper::GlobalLoadStreamB
    GlobalLoadStream< GemmOperand::kB, GlobalLoadIteratorB, SharedStoreIteratorB, GlobalTransformerB > GlobalLoadStreamB
    The stream to load B from global memory to shared memory.
    Definition: hgemm_traits.h:325
    +
    cutlass::gemm::HgemmTraitsHelper::SharedLoadStreamA
    SharedLoadStream< SharedLoadIteratorA > SharedLoadStreamA
    The stream to load A from shared memory.
    Definition: hgemm_traits.h:334
    gemm_epilogue_traits.h
    Defines structural properties of the GEMM epilogue.
    +
    cutlass::gemm::HgemmTraitsHelper::SharedStoreIteratorB
    TileStoreIterator< typename GemmTileTraitsHelperB::SharedStoreTileTraits, typename GemmTileTraitsHelperB::SharedStoreTileTraits::Scalar, IteratorAdvance::kH, MemorySpace::kShared > SharedStoreIteratorB
    The iterator to store B to shared memory.
    Definition: hgemm_traits.h:319
    +
    cutlass::gemm::HgemmTraitsHelper::SharedStoreIteratorA
    TileStoreIterator< typename GemmTileTraitsHelperA::SharedStoreTileTraits, typename GemmTileTraitsHelperA::SharedStoreTileTraits::Scalar, IteratorAdvance::kH, MemorySpace::kShared > SharedStoreIteratorA
    The iterator to store A to shared memory.
    Definition: hgemm_traits.h:300
    cutlass::gemm::HgemmSwizzle
    Definition: hgemm_swizzle.h:40
    convert.h
    Defines conversion operations among Fragments of different base type.
    -
    cutlass::gemm::HgemmTransformerB< MatrixLayout::kRowMajor, Iterator_ >::Transformer
    Convert< typename Iterator_::Fragment, typename Iterator_::Fragment > Transformer
    Definition: hgemm_traits.h:114
    +
    cutlass::gemm::HgemmTransformerB< MatrixLayout::kRowMajor, Iterator_ >::Transformer
    Convert< typename Iterator_::Fragment, typename Iterator_::Fragment > Transformer
    Definition: hgemm_traits.h:120
    cutlass::gemm::HgemmConfig
    Definition: hgemm_traits.h:57
    -
    cutlass::TileStoreIterator
    An iterator implementing Tile Store Iterator Concept for storing a tile to memory.
    Definition: tile_iterator.h:620
    -
    cutlass::gemm::HgemmTraitsHelper::SharedStoreIteratorB
    TileStoreIterator< typename GemmTileTraitsHelperB::SharedStoreTileTraits, typename GemmTileTraitsHelperB::SharedStoreTileTraits::Scalar, IteratorAdvance::kH, MemorySpace::kShared > SharedStoreIteratorB
    The iterator to store B to shared memory.
    Definition: hgemm_traits.h:307
    +
    cutlass::TileStoreIterator
    An iterator implementing Tile Store Iterator Concept for storing a tile to memory.
    Definition: tile_iterator.h:836
    -Generated on Thu May 3 2018 16:36:12 for Cutlass by   +Generated on Wed Sep 19 2018 13:58:50 for Cutlass by   doxygen 1.8.14
    diff --git a/docs/hierarchy.html b/docs/hierarchy.html index 25ba6bda..865698e5 100644 --- a/docs/hierarchy.html +++ b/docs/hierarchy.html @@ -73,7 +73,7 @@ $(function() {
    This inheritance list is sorted roughly, but not completely, alphabetically:
    -
    [detail level 123]

    @@ -135,7 +135,7 @@ Namespaces

    +
    [detail level 1234]
    @@ -94,316 +94,389 @@ $(function() { - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + +
     Ccutlass::platform::aligned_chunk< Align >
     Ccutlass::platform::aligned_storage< Len, Align >Std::aligned_storage
     ▶Ccutlass::AlignedStruct< kAlignment_ >
     Ccutlass::platform::alignment_of< ulong4 >
     Ccutlass::platform::alignment_of< ulonglong2 >
     Ccutlass::platform::alignment_of< ulonglong4 >
     Ccutlass::gemm::ClearAccumulators< Scalar_, kLanes_ >
     Ccutlass::ComputeOffsetFromShape< Shape_ >Compute the offset for the given coordinates in a cube
     Ccutlass::ComputeOffsetFromShape< Shape< 1, kSh_, kSw_, 1 > >Compute the offset for the given coordinates in a cube with one channel and a depth of 1
     Ccutlass::ComputeOffsetFromShape< Shape< 1, kSh_, kSw_, kSc_ > >Compute the offset for the given coordinates in a cube with a depth of 1
     Ccutlass::ComputeOffsetFromStrides< Strides_ >Compute the offset for the given coordinates in a cube
     Ccutlass::ComputeOffsetFromStrides< Shape< 1, S_h_, S_w_, 1 > >Compute the offset for the given coordinates in a cube with one channel and a depth of 1
     Ccutlass::ComputeOffsetFromStrides< Shape< 1, S_h_, S_w_, S_c_ > >Compute the offset for the given coordinates in a cube with a depth of 1
     Ccutlass::ComputeThreadOffsetFromStrides< Threads_, Strides_ >Decompose threadId.x into coordinate of a cube whose dimensions are specified by Threads_. Afterwards compute the offset of those coordinates using Strides_
     Ccutlass::ComputeThreadOffsetFromStrides< Shape< 1, T_h_, T_w_, 1 >, Shape< 1, S_h_, S_w_, 1 > >Specialization for D=1 and C=1
     Ccutlass::ComputeThreadOffsetFromStrides< Shape< 1, T_h_, T_w_, T_c_ >, Shape< 1, S_h_, S_w_, S_c_ > >Specialization for D=1
     Ccutlass::platform::conditional< B, T, F >Std::conditional (true specialization)
     Ccutlass::platform::conditional< false, T, F >Std::conditional (false specialization)
     Ccutlass::PredicateVector< kPredicates_, kPredicatesPerByte_, kPredicateStart_ >::ConstIteratorA const iterator implementing Predicate Iterator Concept enabling sequential read-only access to prediactes
     Ccutlass::ConstPredicateTileAdapter< PredicateVector_, Iterations_ >Adapter to enable random access to predicates via logical coordinate within a tile
     Ccutlass::Convert< InputFragment_, OutputFragment_ >
     Ccutlass::Convert< Fragment< InputScalar_, kScalars_ >, Fragment< OutputScalar_, kScalars_ > >
     Ccutlass::Coord< N_ >Statically-sized array specifying Coords within a tensor
     Ccutlass::Coord< 4 >
     Ccutlass::Coord< Rank >
     Ccutlass::Copy< Fragment_ >
     Ccutlass::platform::default_delete< T >Default deleter
     Ccutlass::platform::default_delete< T[]>Partial specialization for deleting array types
     Ccutlass::divide_assert< Dividend, Divisor >
     Ccutlass::platform::is_base_of_helper< BaseT, DerivedT >::dummy< B, D >
     Ccutlass::platform::enable_if< C, T >Std::enable_if (true specialization)
     Ccutlass::platform::enable_if< false, T >Std::enable_if (false specialization)
     Ccutlass::Extent< T >Returns the extent of a scalar or vector
     Ccutlass::Extent< Vector< T, Lanes > >Returns the number of lanes of a vector if need be
     Ccutlass::Extent< Vector< T, Lanes > const >Returns the number of lanes of a vector if need be
     Ccutlass::FragmentConstIterator< Fragment_, Iterations_, AccessType_ >
     Ccutlass::FragmentIterator< Fragment_, Iterations_, AccessType_ >A template defining Fragment Iterator Concept
     Ccutlass::FragmentLoad< kIteratorFragment, kAccessSize, Scalar_, Memory_, FragmentElement_, kStride >
     Ccutlass::FragmentLoad< IteratorFragment::kScalar, kAccessSize, Scalar_, Memory_, FragmentElement_, kStride >
     Ccutlass::FragmentLoad< IteratorFragment::kWmmaMatrix, kAccessSize, Scalar_, Memory_, FragmentElement_, kStride >
     Ccutlass::gemm::FragmentMultiplyAdd< Scalar_ >
     Ccutlass::gemm::FragmentMultiplyAdd< half >
     Ccutlass::FragmentStore< kIteratorFragment, kAccessSize, Scalar_, Memory_, FragmentElement_, kStride >
     Ccutlass::FragmentStore< IteratorFragment::kScalar, kAccessSize, Scalar_, Memory_, FragmentElement_, kStride >
     Ccutlass::FragmentStore< IteratorFragment::kWmmaMatrix, kAccessSize, Scalar_, Memory_, FragmentElement_, kStride >
     Ccutlass::gemm::Gemm< GemmTraits_ >
     Ccutlass::gemm::GemmConfig< ScalarA_, ScalarB_, ScalarC_, ScalarD_, OutputTile_, MultiplyAdd_, kScalarsPerLdgA_, kScalarsPerStsA_, kScalarsPerLdsA_, kScalarsPerLdgB_, kScalarsPerStsB_, kScalarsPerLdsB_, kScalarsPerLdgCAndStgD_, kScalarsPerStsD_, kScalarsPerLdsD_, kStages_ >
     ▶Ccutlass::gemm::GemmConfig< double, double, double, double, OutputTile_, ThreadMultiplyAdd< AccumulatorsPerThread_, Shape< 1, 4, 8 >, double, double, double >, kScalarsPerLdgA_, kScalarsPerLdgA_, 2, kScalarsPerLdgB_, kScalarsPerLdgB_, 2, 1, 2, 1, 2 >
     ▶Ccutlass::gemm::GemmConfig< float, float, float, float, OutputTile_, ThreadMultiplyAdd< AccumulatorsPerThread_, Shape< 1, 4, 8 >, float, float, float >, kScalarsPerLdgA_, kScalarsPerLdgA_, 4, kScalarsPerLdgB_, kScalarsPerLdgB_, 4, 1, 4, 1, 2 >
     ▶Ccutlass::gemm::GemmConfig< half, half, half, half, OutputTile_, ThreadMultiplyAdd< AccumulatorsPerThread_, Shape< 1, 4, 8 >, half, half, half >, kScalarsPerLdgA_, kScalarsPerLdgA_, 8, kScalarsPerLdgB_, kScalarsPerLdgB_, 8, 2, 8, 2, 2 >
     ▶Ccutlass::gemm::GemmConfig< int8_t, int8_t, int8_t, int8_t, OutputTile_, ThreadMultiplyAdd< AccumulatorsPerThread_, Shape< 1, 4, 8 >, int8_t, int8_t, int >, 4, 4, 16, 4, 4, 16, 4, 4, 4, 2 >
     ▶Ccutlass::gemm::GemmConfig< int8_t, int8_t, ScalarD_, ScalarD_, OutputTile_, ThreadMultiplyAdd< AccumulatorsPerThread_, Shape< 1, 4, 8 >, int8_t, int8_t, int >, 4, 4, 16, 4, 4, 16, 1, 4, 1, 2 >
     Ccutlass::gemm::GemmDesc< Scalar_, Index_ >
     ▶Ccutlass::gemm::GemmEpilogue< GemmEpilogueTraits_ >
     Ccutlass::gemm::GemmEpilogueTraits< OutputTile_, Accumulators_, GlobalLoadIteratorC_, GlobalTransformerC_, GlobalTransformerD_, GlobalStoreIteratorD_, SharedStoreIteratorD_, SharedStoreTransformerD_, SharedLoadIteratorD_, Iterations_, Delta_, Functor_, Index_ >
     ▶Ccutlass::gemm::GemmEpilogueTraits< GemmConfig_::OutputTile, GemmConfig_::Accumulators, Helper_::GlobalLoadIteratorC, Helper_::GlobalTransformerC, Helper_::GlobalTransformerD, Helper_::GlobalStoreIteratorD, Helper_::SharedStoreIteratorD, Helper_::SharedStoreTransformerD, Helper_::SharedLoadIteratorD, Helper_::Iterations, Helper_::Delta, EpilogueFunctor_, Index_ >
     ▶Ccutlass::gemm::GemmEpilogueTraits< IgemmConfig_::OutputTile, IgemmConfig_::Accumulators, Helper_::GlobalLoadIteratorC, Helper_::GlobalTransformerC, Helper_::GlobalTransformerD, Helper_::GlobalStoreIteratorD, Helper_::SharedStoreIteratorD, Helper_::SharedStoreTransformerD, Helper_::SharedLoadIteratorD, Helper_::Iterations, Helper_::Delta, EpilogueFunctor_, Index_ >
     Ccutlass::gemm::GemmEpilogueTraitsHelper< GemmConfig_, EpilogueFunctor_, Index_ >
     ▶Ccutlass::gemm::GemmEpilogueTraitsHelper< IgemmConfig_, EpilogueFunctor_, Index_ >
     ▶Ccutlass::gemm::GemmGlobalTileTraits< kOperand_, kLayout_, Scalar_, Tile_, Threads_, kAccessSize_ >
     ▶Ccutlass::gemm::GemmGlobalTileTraits< GemmOperand::kC, MatrixLayout::kColumnMajor, Scalar_, Tile_, Threads_, kAccessSize_ >
     Ccutlass::gemm::GemmMultiplicandTraits< ThreadBlockTile_, Usage, Layout >
     Ccutlass::GemmOperandGemm operand - D = A * B + C
     Ccutlass::gemm::GemmOperandTraitsAb< kOperand_, kLayout_ >Helper to describe attributes of GEMM matrix operands
     Ccutlass::gemm::GemmSharedLoadTileATraits< Scalar_, OutputTile_, Warps_, ThreadsPerWarp_, InstructionShape_, kStages_, kScalarsPerLds_, kSkew_ >
     Ccutlass::gemm::GemmSharedLoadTileBTraits< Scalar_, OutputTile_, Warps_, ThreadsPerWarp_, InstructionShape_, kStages_, kScalarsPerLds_, kSkew_ >
     Ccutlass::gemm::GemmSharedLoadTileDTraits< Scalar_, OutputTile_, Warps_, ThreadsPerWarp_, kTileH_, kScalarsPerLds_, kSkew_ >
     Ccutlass::gemm::GemmSharedStoreTileAbTraits< Scalar_, Tile_, Threads_, kScalarsPerSts_ >
     Ccutlass::gemm::GemmSharedStoreTileDTraits< Scalar_, OutputTile_, Warps_, ThreadsPerWarp_, kScalarsPerSts_, kSkew_ >
     Ccutlass::gemm::GemmSharedStoreWithSkewTileAbTraits< Scalar_, Tile_, Threads_, kScalarsPerSts_, kSkew_ >
     Ccutlass::gemm::GemmTileTraitsHelperA< Kind, GemmConfig_ >
     ▶Ccutlass::gemm::GemmTileTraitsHelperA< kLayout_, GemmConfig_ >
     ▶Ccutlass::gemm::GemmTileTraitsHelperA< MatrixLayout::kColumnMajor, GemmConfig_ >
     ▶Ccutlass::gemm::GemmTileTraitsHelperA< MatrixLayout::kRowMajor, GemmConfig_ >
     Ccutlass::gemm::GemmTileTraitsHelperB< Kind, GemmConfig_ >
     ▶Ccutlass::gemm::GemmTileTraitsHelperB< kLayout_, GemmConfig_ >
     ▶Ccutlass::gemm::GemmTileTraitsHelperB< MatrixLayout::kColumnMajor, GemmConfig_ >
     ▶Ccutlass::gemm::GemmTileTraitsHelperB< MatrixLayout::kRowMajor, GemmConfig_ >
     Ccutlass::gemm::GemmTraits< GemmConfig_, GlobalLoadStreamA_, GlobalLoadStreamB_, SharedLoadStreamA_, SharedLoadStreamB_, Epilogue_, BlockSwizzle_, Index_, ClearAccumulators_ >
     ▶Ccutlass::gemm::GemmTraits< GemmConfig_, Helper_::GlobalLoadStreamA, Helper_::GlobalLoadStreamB, Helper_::SharedLoadStreamA, Helper_::SharedLoadStreamB, Epilogue_, IdentityBlockSwizzle, Index_, ClearAccumulators< GemmConfig_::Accumulators::Element > >
     ▶Ccutlass::gemm::GemmTraits< GemmConfig_, SimplifiedGemmTraitsHelper< GemmTileTraitsHelperA< kLayoutA_, GemmConfig_ >, GemmTileTraitsHelperB< kLayoutB_, GemmConfig_ >, Index_ > ::GlobalLoadStreamA, SimplifiedGemmTraitsHelper< GemmTileTraitsHelperA< kLayoutA_, GemmConfig_ >, GemmTileTraitsHelperB< kLayoutB_, GemmConfig_ >, Index_ > ::GlobalLoadStreamB, SimplifiedGemmTraitsHelper< GemmTileTraitsHelperA< kLayoutA_, GemmConfig_ >, GemmTileTraitsHelperB< kLayoutB_, GemmConfig_ >, Index_ > ::SharedLoadStreamA, SimplifiedGemmTraitsHelper< GemmTileTraitsHelperA< kLayoutA_, GemmConfig_ >, GemmTileTraitsHelperB< kLayoutB_, GemmConfig_ >, Index_ > ::SharedLoadStreamB, GemmEpilogue< GemmEpilogueTraits_ >, IdentityBlockSwizzle, Index_, ClearAccumulators< GemmConfig_::Accumulators::Element > >
     ▶Ccutlass::gemm::GemmTraits< Helper_::GemmConfig, Helper_::GlobalLoadStreamA, Helper_::GlobalLoadStreamB, Helper_::SharedLoadStreamA, Helper_::SharedLoadStreamB, Helper_::Epilogue, IdentityBlockSwizzle, Index_, Helper_::ClearAccumulators >
     Ccutlass::gemm::GetExtent< kOperand_, Tile_ >
     Ccutlass::gemm::GetExtent< GemmOperand::kA, Tile_ >
     Ccutlass::gemm::GetExtent< GemmOperand::kB, Tile_ >
     Ccutlass::gemm::GemmTraits< GemmConfig_, GlobalLoadStreamA_, GlobalLoadStreamB_, SharedLoadStreamA_, SharedLoadStreamB_, Epilogue_, BlockSwizzle_, Index_, ClearAccumulators_ >::GlobalLoadStreamAssemble the global load streams for A/B
     ▶Ccutlass::gemm::GlobalLoadStreamBase< LoadIterator_, StoreIterator_, Transformer_ >
     Ccutlass::platform::greater< T >Std::greater
     Ccutlass::gemm::HgemmSwizzle< GlobalIterator_ >
     Ccutlass::gemm::HgemmTraitsHelper< kLayoutA_, kLayoutB_, OutputTile_, EpilogueFunctor_, AccumulatorsPerThread_, kScalarsPerLdgA_, kScalarsPerLdgB_, Index_ >
     Ccutlass::gemm::HgemmTransformerA< kLayout_, Iterator_ >
     Ccutlass::gemm::HgemmTransformerA< MatrixLayout::kColumnMajor, Iterator_ >
     Ccutlass::gemm::HgemmTransformerA< MatrixLayout::kRowMajor, Iterator_ >
     Ccutlass::gemm::HgemmTransformerB< kLayout_, Iterator_ >
     Ccutlass::gemm::HgemmTransformerB< MatrixLayout::kColumnMajor, Iterator_ >
     Ccutlass::gemm::HgemmTransformerB< MatrixLayout::kRowMajor, Iterator_ >
     Ccutlass::IdentityDescribes identity elements
     Ccutlass::gemm::IdentityBlockSwizzle
     Ccutlass::gemm::IgemmEpilogueScalar< ScalarD_ >
     Ccutlass::gemm::IgemmEpilogueScalar< int >
     Ccutlass::gemm::IgemmFloatToInt8Converter< kElements_ >
     Ccutlass::gemm::IgemmGlobalLoadTransformer< InputFragment_, OutputScalar_ >
     Ccutlass::gemm::IgemmGlobalLoadTransformer< Fragment< int8_t, kElements_ >, float >
     Ccutlass::gemm::IgemmGlobalStoreTransformer< InputScalar_, OutputFragment_ >
     Ccutlass::gemm::IgemmGlobalStoreTransformer< float, Fragment< int8_t, kElements_ > >
     Ccutlass::gemm::IgemmInt8ToFloatConverter< kElements_ >
     Ccutlass::gemm::IgemmSharedStoreTransformer< InputScalar_, OutputFragment_ >
     Ccutlass::gemm::IgemmSwizzle< GlobalIterator_ >
     Ccutlass::gemm::IgemmTraitsHelper< kLayoutA_, kLayoutB_, OutputTile_, ScalarD_, EpilogueFunctor_, AccumulatorsPerThread_, Index_ >
     Ccutlass::gemm::IgemmTransformerA< kLayout_, Iterator_ >
     Ccutlass::gemm::IgemmTransformerA< MatrixLayout::kColumnMajor, Iterator_ >
     Ccutlass::gemm::IgemmTransformerA< MatrixLayout::kRowMajor, Iterator_ >
     Ccutlass::gemm::IgemmTransformerB< kLayout_, Iterator_ >
     Ccutlass::gemm::IgemmTransformerB< MatrixLayout::kColumnMajor, Iterator_ >
     Ccutlass::gemm::IgemmTransformerB< MatrixLayout::kRowMajor, Iterator_ >
     ▶Ccutlass::platform::integral_constant< value_t, V >Std::integral_constant
     ▶Ccutlass::platform::integral_constant< bool, V >
     ▶Ccutlass::platform::integral_constant< bool,(is_arithmetic< T >::value||is_void< T >::value||is_same< nullptr_t, remove_cv< T >::type >::value)>
     ▶Ccutlass::platform::integral_constant< bool,(is_base_of_helper< remove_cv< BaseT >::type, remove_cv< DerivedT >::type >::value)||(is_same< remove_cv< BaseT >::type, remove_cv< DerivedT >::type >::value)>
     ▶Ccutlass::platform::integral_constant< bool,(is_fundamental< T >::value||is_pointer< T >::value)>
     ▶Ccutlass::platform::integral_constant< bool,(is_integral< T >::value||is_floating_point< T >::value)>
     ▶Ccutlass::platform::integral_constant< bool,(is_same< float, remove_cv< T >::type >::value||is_same< double, remove_cv< T >::type >::value)>
     ▶Ccutlass::platform::integral_constant< bool,(N &(N - 1))==0 >
     Ccutlass::platform::is_base_of_helper< BaseT, DerivedT >Helper for std::is_base_of
     Ccutlass::PredicateVector< kPredicates_, kPredicatesPerByte_, kPredicateStart_ >::IteratorAn iterator implementing Predicate Iterator Concept enabling sequential read and write access to predicates
     Ccutlass::IteratorAdvanceSpecifies dimension in which post-increment accesses advance
     Ccutlass::IteratorFragmentSpecifies whether iterator storage fragment consists of Scalar values or WMMA matrix
     Ccutlass::platform::less< T >Std::less
     Ccutlass::gemm::LinearScaling< Scalar_, FragmentMultiplyAdd_ >Functor to compute linear combination of fragments
     Ccutlass::Load< Scalar_, Lanes_, Memory_, bool, size_t >
     Ccutlass::Load< double, 2, Memory_, true, 16 >
     Ccutlass::Load< Scalar_, Lanes_, Memory_, true, 16 >
     Ccutlass::Load< Scalar_, Lanes_, Memory_, true, 4 >
     Ccutlass::Load< Scalar_, Lanes_, Memory_, true, 8 >
     Ccutlass::log2_down< N, CurrentVal, Count >
     Ccutlass::log2_down< N, 1, Count >
     Ccutlass::log2_up< N, CurrentVal, Count >
     Ccutlass::log2_up< N, 1, Count >
     Ccutlass::gemm::GemmTraits< GemmConfig_, GlobalLoadStreamA_, GlobalLoadStreamB_, SharedLoadStreamA_, SharedLoadStreamB_, Epilogue_, BlockSwizzle_, Index_, ClearAccumulators_ >::MainLoopSharedStorage
     Ccutlass::MatrixLayoutDescribes layouts of matrices
     Ccutlass::MemorySpaceEnum to specify which memory space data resides in
     Ccutlass::platform::nullptr_tStd::nullptr_t
     Ccutlass::platform::alignment_of< value_t >::pad
     Ccutlass::gemm::WmmaGemmGlobalIteratorCd< TileTraits_, Index_ >::ParamsThe params
     ▶CParams
     Ccutlass::gemm::GemmTraits< GemmConfig_, GlobalLoadStreamA_, GlobalLoadStreamB_, SharedLoadStreamA_, SharedLoadStreamB_, Epilogue_, BlockSwizzle_, Index_, ClearAccumulators_ >::ParamsThe params
     Ccutlass::gemm::GlobalLoadStreamBase< LoadIterator_, StoreIterator_, Transformer_ >::ParamsThe params
     ▶Ccutlass::TileIteratorBase< Traits_, Scalar_, Advance_, MemorySpace, Index_, FragmentElement_, IteratorFragment_, Skew_ >::ParamsParameters to the iterator
     Ccutlass::gemm::GemmGlobalIteratorCd< TileTraits_, Index_ >::ParamsThe params
     Ccutlass::gemm::GemmEpilogueTraits< OutputTile_, Accumulators_, GlobalLoadIteratorC_, GlobalTransformerC_, GlobalTransformerD_, GlobalStoreIteratorD_, SharedStoreIteratorD_, SharedStoreTransformerD_, SharedLoadIteratorD_, Iterations_, Delta_, Functor_, Index_ >::ParamsThe params
     Ccutlass::gemm::SharedLoadStream< Iterator_, Transformer_ >::ParamsThe params
     Ccutlass::gemm::LinearScaling< Scalar_, FragmentMultiplyAdd_ >::ParamsThe parameters
     Ccutlass::platform::plus< T >Platform::plus
     Ccutlass::PredicateTileAdapter< PredicateVector_, Iterations_ >Adapter to enable random access to predicates via logical coordinate within a tile
     Ccutlass::PredicateVector< kPredicates_, kPredicatesPerByte_, kPredicateStart_ >Statically sized array of bits implementing
     Ccutlass::PredicateVector< Base::Iterations::kW >
     Ccutlass::PredicateVector< ShapeCount< typename Base::Iterations >::kCount >
     Ccutlass::gemm::ProjectOperand< operand, Kstrided >
     Ccutlass::gemm::ProjectOperand< GemmOperand::kA, Kstrided >Project A operand - (0, K, M)
     Ccutlass::gemm::ProjectOperand< GemmOperand::kB, Kstrided >Project B operand - (0, K, N)
     Ccutlass::gemm::ProjectOperand< GemmOperand::kC, true >Project C operand - (0, N, M)
     Ccutlass::gemm::ProjectOperand< GemmOperand::kD, true >Project D operand - (0, N, M)
     Ccutlass::platform::remove_const< T >Std::remove_const (non-const specialization)
     Ccutlass::platform::remove_const< const T >Std::remove_const (const specialization)
     Ccutlass::platform::remove_cv< T >Std::remove_cv
     Ccutlass::platform::remove_volatile< T >Std::remove_volatile (non-volatile specialization)
     Ccutlass::platform::remove_volatile< volatile T >Std::remove_volatile (volatile specialization)
     Ccutlass::gemm::ReshapeThreads< Tile_, Threads_, bool >
     Ccutlass::gemm::ReshapeThreads< Tile_, Threads_, true >
     Ccutlass::ReshapeTile< Tile_, kAccessSize_, bool >
     Ccutlass::ReshapeTile< Tile_, kAccessSize_, true >
     Ccutlass::Shape< kD_, kH_, kW_, kC_ >A Shape implementing Layout Concept describing the dimensions of a cube
     Ccutlass::ShapeAdd< A_, B_ >
     Ccutlass::ShapeCount< Shape >Compute derived counted of a Layout Concept based class
     Ccutlass::ShapeDiv< A_, B_ >
     Ccutlass::ShapeMax< A_, B_ >
     Ccutlass::ShapeMin< A_, B_ >
     Ccutlass::ShapeMul< A_, B_ >
     Ccutlass::ShapeScale< A_, kScale_ >
     Ccutlass::ShapeStrides< Shape_ >
     Ccutlass::ShapeSub< A_, B_ >
     Ccutlass::gemm::GemmTraits< GemmConfig_, GlobalLoadStreamA_, GlobalLoadStreamB_, SharedLoadStreamA_, SharedLoadStreamB_, Epilogue_, BlockSwizzle_, Index_, ClearAccumulators_ >::SharedLoadStreamAssemble the shared load stream for A/B
     Ccutlass::gemm::SharedLoadStream< Iterator_, Transformer_ >
     Ccutlass::gemm::ClearAccumulators< Scalar_, kLanes_ >::SharedStorageThe shared storage
     Ccutlass::gemm::GemmEpilogueTraits< OutputTile_, Accumulators_, GlobalLoadIteratorC_, GlobalTransformerC_, GlobalTransformerD_, GlobalStoreIteratorD_, SharedStoreIteratorD_, SharedStoreTransformerD_, SharedLoadIteratorD_, Iterations_, Delta_, Functor_, Index_ >::SharedStorageThe shared memory to swizzle the data in the epilogue
     Ccutlass::gemm::GemmTraits< GemmConfig_, GlobalLoadStreamA_, GlobalLoadStreamB_, SharedLoadStreamA_, SharedLoadStreamB_, Epilogue_, BlockSwizzle_, Index_, ClearAccumulators_ >::SharedStorageThe storage in shared memory
     Ccutlass::gemm::GlobalLoadStreamBase< LoadIterator_, StoreIterator_, Transformer_ >::SharedStorageThe storage in shared memory needed by that stream
     Ccutlass::gemm::SimplifiedGemmTraitsHelper< GemmTileTraitsHelperA_, GemmTileTraitsHelperB_, Index_ >
     Ccutlass::sqrt_est< N >
     Ccutlass::StorageType< kAlignment_ >
     Ccutlass::StorageType< 1 >
     Ccutlass::StorageType< 2 >
     Ccutlass::StorageType< 4 >
     Ccutlass::Store< Scalar_, Lanes_, Memory_, bool, size_t >
     Ccutlass::Store< double, 2, Memory_, true, 16 >
     Ccutlass::Store< Scalar_, Lanes_, Memory_, true, 16 >
     Ccutlass::Store< Scalar_, Lanes_, Memory_, true, 4 >
     Ccutlass::Store< Scalar_, Lanes_, Memory_, true, 8 >
     Ccutlass::gemm::GemmTraits< GemmConfig_, GlobalLoadStreamA_, GlobalLoadStreamB_, SharedLoadStreamA_, SharedLoadStreamB_, Epilogue_, BlockSwizzle_, Index_, ClearAccumulators_ >::StreamSharedStorage< GlobalLoadStream_, SharedLoadStream_ >
     Ccutlass::gemm::GemmEpilogueTraits< OutputTile_, Accumulators_, GlobalLoadIteratorC_, GlobalTransformerC_, GlobalTransformerD_, GlobalStoreIteratorD_, SharedStoreIteratorD_, SharedStoreTransformerD_, SharedLoadIteratorD_, Iterations_, Delta_, Functor_, Index_ >::StreamSharedStorageThe shared memory storage to exchange data
     Ccutlass::gemm::GemmTraits< GemmConfig_, GlobalLoadStreamA_, GlobalLoadStreamB_, SharedLoadStreamA_, SharedLoadStreamB_, Epilogue_, BlockSwizzle_, Index_, ClearAccumulators_ >::StreamSharedStorage< GlobalLoadStreamA, SharedLoadStreamA >
     Ccutlass::gemm::GemmTraits< GemmConfig_, GlobalLoadStreamA_, GlobalLoadStreamB_, SharedLoadStreamA_, SharedLoadStreamB_, Epilogue_, BlockSwizzle_, Index_, ClearAccumulators_ >::StreamSharedStorage< GlobalLoadStreamB, SharedLoadStreamB >
     Ccutlass::TensorRef< Storage_, Rank_ >Structure modeling a pointer and stride into a tensor
     ▶Ccutlass::TensorRef< T, 4 >
     Ccutlass::gemm::ThreadMultiplyAdd< AccumulatorsPerThread_, ThreadsPerWarp_, ScalarA_, ScalarB_, ScalarC_ >Template performing matrix multiply-add operation within a thread
     Ccutlass::gemm::ThreadMultiplyAdd< AccumulatorsPerThread_, ThreadsPerWarp_, half, half, half >Template performing matrix multiply-add operation within a thread
     Ccutlass::gemm::ThreadMultiplyAdd< AccumulatorsPerThread_, ThreadsPerWarp_, int8_t, int8_t, int >Template performing matrix multiply-add operation within a thread
     Ccutlass::gemm::GemmSharedLoadTileBTraits< Scalar_, OutputTile_, Warps_, ThreadsPerWarp_, InstructionShape_, kStages_, kScalarsPerLds_, kSkew_ >::ThreadOffsetComputes the thread offset in (H, W) based on thread ID
     Ccutlass::gemm::GemmGlobalTileCdTraits< Scalar_, Tile_, Threads_, kStrideH_, kAccessSize_ >::ThreadOffsetComputes the thread offset in (H, W) based on thread ID
     Ccutlass::gemm::IgemmContiguousGlobalTileTraits< kOperand_, kLayout_, Scalar_, Tile_, Threads_, kAccessSize_ >::ThreadOffsetComputes the thread offset in (H, W) based on thread ID
     Ccutlass::gemm::GemmGlobalTileTraits< kOperand_, kLayout_, Scalar_, Tile_, Threads_, kAccessSize_ >::ThreadOffsetComputes the thread offset in (H, W) based on thread ID
     Ccutlass::gemm::GemmSharedLoadTileDTraits< Scalar_, OutputTile_, Warps_, ThreadsPerWarp_, kTileH_, kScalarsPerLds_, kSkew_ >::ThreadOffsetComputes the thread offset in (H, W) based on thread ID
     Ccutlass::gemm::GemmSharedLoadTileATraits< Scalar_, OutputTile_, Warps_, ThreadsPerWarp_, InstructionShape_, kStages_, kScalarsPerLds_, kSkew_ >::ThreadOffsetComputes the thread offset in (H, W) based on thread ID
     Ccutlass::gemm::GemmSharedStoreTileDTraits< Scalar_, OutputTile_, Warps_, ThreadsPerWarp_, kScalarsPerSts_, kSkew_ >::ThreadOffsetComputes the thread offset in (H, W) based on thread ID
     Ccutlass::gemm::HgemmCrosswiseGlobalTileTraits< kOperand_, kLayout_, Scalar_, Tile_, Threads_, kAccessSize_ >::ThreadOffsetComputes the thread offset in (H, W) based on thread ID
     Ccutlass::gemm::GemmSharedStoreTileAbTraits< Scalar_, Tile_, Threads_, kScalarsPerSts_ >::ThreadOffset
     Ccutlass::TileTraitsWarpRake< Tile_, Threads >::ThreadOffsetComputes the thread offset in (H, W) based on thread ID
     Ccutlass::gemm::GemmSharedStoreWithSkewTileAbTraits< Scalar_, Tile_, Threads_, kScalarsPerSts_, kSkew_ >::ThreadOffset
     Ccutlass::gemm::WmmaGemmGlobalIteratorCdTraits< Scalar_, Tile_, Threads_, kAccessSize_ >::ThreadOffsetComputes the thread offset in (H, W) based on thread ID
     Ccutlass::TiledThreadOffset< ThreadShape >Basic thread offset function computed from a thread shape
     ▶Ccutlass::TileIteratorBase< Traits_, Scalar_, Advance_, MemorySpace, Index_, FragmentElement_, IteratorFragment_, Skew_ >Iterator for accessing a stripmined tile in memory
     ▶Ccutlass::TileIteratorBase< TileTraits_, TileTraits_::Scalar, Advance_, MemorySpace, Index_, TileTraits_::Scalar, IteratorFragment::kScalar, Shape< 0, 0, 0, 0 > >
     ▶Ccutlass::TileIteratorBase< TileTraits_, TileTraits_::Scalar, IteratorAdvance::kH, MemorySpace::kGlobal, Index_ >
     Ccutlass::TileTraits< Tile_, Delta_, Iterations_, ThreadOffset_ >A template defining Tile Traits Concept
     Ccutlass::TileTraitsContiguousMajor< Tile_, Threads >
     Ccutlass::TileTraitsStandard< Tile_, Threads >Chooses 'best' shape to enable warp raking along contiguous dimension if possible
     Ccutlass::TileTraitsStrideMajor< Tile_, Threads >
     Ccutlass::TileTraitsWarpRake< Tile_, Threads >Tiling in which warps rake across the contiguous dimension
     Ccutlass::PredicateVector< kPredicates_, kPredicatesPerByte_, kPredicateStart_ >::TrivialIteratorIterator that always returns true
     Ccutlass::TrivialPredicateTileAdapterAlways returns true predicate
     Ccutlass::platform::unique_ptr< T, Deleter >Std::unique_ptr
     Ccutlass::Vector< Scalar_, kLanes_ >
     Ccutlass::Vector< half, kLanes_ >
     Ccutlass::Vectorize< Element_, kLanes_ >
     Ccutlass::Vectorize< Element_, 1 >
     Ccutlass::VectorTraits< T >Traits describing properties of vectors and scalar-as-vectors
     Ccutlass::VectorTraits< Vector< T, Lanes > >Partial specialization for actual cutlass::Vector
     Ccutlass::VectorTraits< Vector< T, Lanes > const >Partial specialization for actual cutlass::Vector
     Ccutlass::bin1_t
     Ccutlass::gemm::ClearAccumulators< Scalar_, kLanes_ >
     Ccutlass::MatrixLayout::ColumnMajorMapping function for column-major matrices
     Ccutlass::MatrixLayout::ColumnMajorBlockLinear< BlockRows, BlockColumns >
     Ccutlass::gemm::ColumnMajorBlockSwizzle< groupCols, swDirection >
     Ccutlass::MatrixLayout::ColumnMajorInterleaved< Interleave >
     Ccutlass::platform::complex< T >
     Ccutlass::ComputeOffsetFromShape< Shape_ >Compute the offset for the given coordinates in a cube
     Ccutlass::ComputeOffsetFromStrides< Strides_ >Compute the offset for the given coordinates in a cube
     Ccutlass::ComputeThreadOffsetFromStrides< Threads_, Strides_ >Decompose threadId.x into coordinate of a cube whose dimensions are specified by Threads_. Afterwards compute the offset of those coordinates using Strides_
     Ccutlass::ComputeThreadOffsetFromStrides< Shape< 1, T_h_, T_w_, 1 >, Shape< 1, S_h_, S_w_, 1 > >Specialization for D=1 and C=1
     Ccutlass::ComputeThreadOffsetFromStrides< Shape< 1, T_h_, T_w_, T_c_ >, Shape< 1, S_h_, S_w_, S_c_ > >Specialization for D=1
     Ccutlass::platform::conditional< B, T, F >Std::conditional (true specialization)
     Ccutlass::platform::conditional< false, T, F >Std::conditional (false specialization)
     Ccutlass::PredicateVector< kPredicates_, kPredicatesPerByte_, kPredicateStart_ >::ConstIteratorA const iterator implementing Predicate Iterator Concept enabling sequential read-only access to prediactes
     Ccutlass::TensorRefBatchStrided< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::ConstIteratorConstant iterator over tensors implied by TensorRefBatchStrided
     Ccutlass::TensorRefArray< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >::ConstIteratorTensorRefIterator over TensorRef objects in TensorRefArray
     Ccutlass::ConstPredicateTileAdapter< PredicateVector_, Iterations_ >Adapter to enable random access to predicates via logical coordinate within a tile
     Ccutlass::MatrixLayout::ContiguousLayout
     Ccutlass::Convert< InputFragment_, OutputFragment_ >
     Ccutlass::Convert< Fragment< InputScalar_, kScalars_ >, Fragment< OutputScalar_, kScalars_ > >
     Ccutlass::Coord< Rank_, Index_ >Statically-sized array specifying Coords within a tensor
     ▶Ccutlass::Coord< 2, int >
     Ccutlass::Coord< 3 >
     Ccutlass::Coord< 4 >
     ▶Ccutlass::Coord< 4, Index_ >
     ▶Ccutlass::Coord< 4, int >
     Ccutlass::Coord< kStorageRank - 1 >
     Ccutlass::Copy< Fragment_ >
     CDebugType< T >
     CDebugValue< Value >
     Ccutlass::platform::default_delete< T >Default deleter
     Ccutlass::platform::default_delete< T[]>Partial specialization for deleting array types
     Ccutlass::divide_assert< Dividend, Divisor >
     Ccutlass::platform::is_base_of_helper< BaseT, DerivedT >::dummy< B, D >
     Ccutlass::DumpType< T >
     Ccutlass::platform::enable_if< C, T >Std::enable_if (true specialization)
     Ccutlass::platform::enable_if< false, T >Std::enable_if (false specialization)
     Ccutlass::Extent< T >Returns the extent of a scalar or vector
     Ccutlass::Extent< Vector< T, Lanes > >Returns the number of lanes of a vector if need be
     Ccutlass::Extent< Vector< T, Lanes > const >Returns the number of lanes of a vector if need be
     Ccutlass::FragmentConstIterator< Fragment_, Iterations_, AccessType_ >
     Ccutlass::FragmentElementTypeSpecifies whether iterator storage fragment consists of Scalar values or WMMA matrix
     Ccutlass::FragmentIterator< Fragment_, Iterations_, AccessType_ >A template defining Fragment Iterator Concept
     Ccutlass::gemm::FragmentMultiplyAdd< ScalarAlphaBeta_, ScalarAccum_, fragMul2 >
     Ccutlass::gemm::FragmentMultiplyAdd< half, half, true >
     Ccutlass::gemm::Gemm< GemmTraits_ >
     Ccutlass::gemm::GemmConfig< ScalarA_, ScalarB_, ScalarC_, ScalarD_, OutputTile_, MultiplyAdd_, kScalarsPerLdgA_, kScalarsPerStsA_, kScalarsPerLdsA_, kScalarsPerLdgB_, kScalarsPerStsB_, kScalarsPerLdsB_, kScalarsPerLdgCAndStgD_, kScalarsPerStsD_, kScalarsPerLdsD_, kStages_, kResidueSeparate_, kResidueInProlog_, kLaunchBounds_ >
     ▶Ccutlass::gemm::GemmConfig< double, double, double, double, OutputTile_, ThreadMultiplyAdd< ThreadGemmShape_, Shape< 1, 4, 8 >, double, double, double >, kScalarsPerLdgA_, kScalarsPerLdgA_, 2, kScalarsPerLdgB_, kScalarsPerLdgB_, 2, 1, 2, 1, 2, false, false, false >
     ▶Ccutlass::gemm::GemmConfig< float, float, float, float, OutputTile_, ThreadMultiplyAdd< ThreadGemmShape_, Shape< 1, 4, 8 >, float, float, float >, kScalarsPerLdgA_, kScalarsPerLdgA_, 4, kScalarsPerLdgB_, kScalarsPerLdgB_, 4, 1, 4, 1, 2, false, true, kLaunchBounds >
     ▶Ccutlass::gemm::GemmConfig< half, half, half, half, OutputTile_, ThreadMultiplyAdd< ThreadGemmShape_, Shape< 1, 4, 8 >, half, half, half >, kScalarsPerLdgA_, kScalarsPerLdgA_, 8, kScalarsPerLdgB_, kScalarsPerLdgB_, 8, 2, 8, 2, 2, false, true, false >
     ▶Ccutlass::gemm::GemmConfig< int8_t, int8_t, int8_t, int8_t, OutputTile_, ThreadMultiplyAdd< ThreadGemmShape_, Shape< 1, 4, 8 >, int8_t, int8_t, int >, 4, 4, 16, 4, 4, 16, 4, 4, 4, 2, false, true, false >
     ▶Ccutlass::gemm::GemmConfig< int8_t, int8_t, ScalarD_, ScalarD_, OutputTile_, ThreadMultiplyAdd< ThreadGemmShape_, Shape< 1, 4, 8 >, int8_t, int8_t, int >, 4, 4, 16, 4, 4, 16, 1, 4, 1, 2, false, false, false >
     ▶Ccutlass::gemm::GemmConfig< ScalarA_, ScalarB_, ScalarC_, ScalarD_, OutputTile_, ThreadMultiplyAdd< ThreadGemmShape_, Shape< 1, 4, 8 >, ScalarA_, ScalarB_, float >, kScalarsPerLdgA_, kScalarsPerLdgA_, 4, kScalarsPerLdgB_, kScalarsPerLdgB_, 4, 1, 4, 1, 2 >
     Ccutlass::gemm::GemmDesc< AType_, BType_, CType_, DType_, SType_, Index_ >GEMM problem description
     ▶Ccutlass::gemm::GemmEpilogue< GemmEpilogueTraits_ >
     Ccutlass::gemm::GemmEpilogueTraits< OutputTile_, Accumulators_, GlobalLoadIteratorC_, GlobalTransformerC_, GlobalTransformerD_, GlobalStoreIteratorD_, SharedStoreIteratorD_, SharedStoreTransformerD_, SharedLoadStreamD_, Iterations_, Delta_, Functor_, Index_ >
     ▶Ccutlass::gemm::GemmEpilogueTraits< GemmConfig_::OutputTile, GemmConfig_::Accumulators, Helper_::GlobalLoadIteratorC, Helper_::GlobalTransformerC, Helper_::GlobalTransformerD, Helper_::GlobalStoreIteratorD, Helper_::SharedStoreIteratorD, Helper_::SharedStoreTransformerD, Helper_::SharedLoadStreamD, Helper_::Iterations, Helper_::Delta, EpilogueFunctor_, Index_ >
     ▶Ccutlass::gemm::GemmEpilogueTraits< IgemmConfig_::OutputTile, IgemmConfig_::Accumulators, Helper_::GlobalLoadIteratorC, Helper_::GlobalTransformerC, Helper_::GlobalTransformerD, Helper_::GlobalStoreIteratorD, Helper_::SharedStoreIteratorD, Helper_::SharedStoreTransformerD, Helper_::SharedLoadStreamD, Helper_::Iterations, Helper_::Delta, EpilogueFunctor_, Index_ >
     Ccutlass::gemm::GemmEpilogueTraitsHelper< GemmConfig_, EpilogueFunctor_, Index_ >
     ▶Ccutlass::gemm::GemmEpilogueTraitsHelper< IgemmConfig_, EpilogueFunctor_, Index_ >
     ▶Ccutlass::gemm::GemmGlobalTileTraits< kOperand_, kLayout_, Scalar_, Tile_, Threads_, kAccessSize_ >
     ▶Ccutlass::gemm::GemmGlobalTileTraits< GemmOperand::kC, MatrixLayout::kColumnMajor, Scalar_, Tile_, Threads_, kAccessSize_ >
     Ccutlass::gemm::GemmMultiplicandTraits< ThreadBlockTile_, Usage, Layout >
     Ccutlass::GemmOperandGemm operand - D = A * B + C
     Ccutlass::gemm::GemmOperandTraitsAb< kOperand_, kLayout_ >Helper to describe attributes of GEMM matrix operands
     Ccutlass::gemm::GemmSharedLoadTileATraits< Scalar_, OutputTile_, Warps_, ThreadsPerWarp_, InstructionShape_, kStages_, kScalarsPerLds_, kSkew_ >
     Ccutlass::gemm::GemmSharedLoadTileBTraits< Scalar_, OutputTile_, Warps_, ThreadsPerWarp_, InstructionShape_, kStages_, kScalarsPerLds_, kSkew_ >
     Ccutlass::gemm::GemmSharedLoadTileDTraits< Scalar_, OutputTile_, Warps_, ThreadsPerWarp_, kTileH_, kScalarsPerLds_, kSkew_ >
     Ccutlass::gemm::GemmSharedStoreTileAbTraits< Scalar_, Tile_, Threads_, kScalarsPerSts_ >
     Ccutlass::gemm::GemmSharedStoreTileDTraits< Scalar_, OutputTile_, Warps_, ThreadsPerWarp_, kScalarsPerSts_, kSkew_ >
     Ccutlass::gemm::GemmSharedStoreWithSkewTileAbTraits< Scalar_, Tile_, Threads_, kScalarsPerSts_, kSkew_ >
     Ccutlass::gemm::GemmTileTraitsHelperA< Kind, GemmConfig_ >
     ▶Ccutlass::gemm::GemmTileTraitsHelperA< kLayout_, GemmConfig_ >
     ▶Ccutlass::gemm::GemmTileTraitsHelperA< MatrixLayout::kColumnMajor, GemmConfig_ >
     ▶Ccutlass::gemm::GemmTileTraitsHelperA< MatrixLayout::kRowMajor, GemmConfig_ >
     Ccutlass::gemm::GemmTileTraitsHelperB< Kind, GemmConfig_ >
     ▶Ccutlass::gemm::GemmTileTraitsHelperB< kLayout_, GemmConfig_ >
     ▶Ccutlass::gemm::GemmTileTraitsHelperB< MatrixLayout::kColumnMajor, GemmConfig_ >
     ▶Ccutlass::gemm::GemmTileTraitsHelperB< MatrixLayout::kRowMajor, GemmConfig_ >
     Ccutlass::gemm::GemmTraits< GemmConfig_, GlobalLoadStreamA_, GlobalLoadStreamB_, SharedLoadStreamA_, SharedLoadStreamB_, Epilogue_, BlockSwizzle_, Index_, ClearAccumulators_ >
     ▶Ccutlass::gemm::GemmTraits< GemmConfig_, Helper_::GlobalLoadStreamA, Helper_::GlobalLoadStreamB, Helper_::SharedLoadStreamA, Helper_::SharedLoadStreamB, Epilogue_, IdentityBlockSwizzle, Index_, ClearAccumulators< GemmConfig_::Accumulators::Element > >
     ▶Ccutlass::gemm::GemmTraits< GemmConfig_, SimplifiedGemmTraitsHelper< GemmTileTraitsHelperA< kLayoutA_, GemmConfig_ >, GemmTileTraitsHelperB< kLayoutB_, GemmConfig_ >, Index_ > ::GlobalLoadStreamA, SimplifiedGemmTraitsHelper< GemmTileTraitsHelperA< kLayoutA_, GemmConfig_ >, GemmTileTraitsHelperB< kLayoutB_, GemmConfig_ >, Index_ > ::GlobalLoadStreamB, SimplifiedGemmTraitsHelper< GemmTileTraitsHelperA< kLayoutA_, GemmConfig_ >, GemmTileTraitsHelperB< kLayoutB_, GemmConfig_ >, Index_ > ::SharedLoadStreamA, SimplifiedGemmTraitsHelper< GemmTileTraitsHelperA< kLayoutA_, GemmConfig_ >, GemmTileTraitsHelperB< kLayoutB_, GemmConfig_ >, Index_ > ::SharedLoadStreamB, GemmEpilogue< GemmEpilogueTraits_ >, IdentityBlockSwizzle, Index_, ClearAccumulators< GemmConfig_::Accumulators::Element > >
     ▶Ccutlass::gemm::GemmTraits< Helper_::GemmConfig, Helper_::GlobalLoadStreamA, Helper_::GlobalLoadStreamB, Helper_::SharedLoadStreamA, Helper_::SharedLoadStreamB, Helper_::Epilogue, IdentityBlockSwizzle, Index_, Helper_::ClearAccumulators >
     Ccutlass::gemm::GetExtent< kOperand_, Tile_ >
     Ccutlass::gemm::GetExtent< GemmOperand::kA, Tile_ >
     Ccutlass::gemm::GetExtent< GemmOperand::kB, Tile_ >
     Ccutlass::gemm::GlobalLoadStream< Operand, LoadIterator_, StoreIterator_, Transformer_ >
     Ccutlass::gemm::GlobalLoadStreamPair< StreamA_, StreamB_, kResidueInProlog_ >Collect the global load streams for multiplicands
     Ccutlass::platform::greater< T >Std::greater
     Ccutlass::gemm::HgemmSwizzle< GlobalIterator_ >
     Ccutlass::gemm::HgemmTraitsHelper< kLayoutA_, kLayoutB_, OutputTile_, EpilogueFunctor_, ThreadGemmShape_, kScalarsPerLdgA_, kScalarsPerLdgB_, Index_ >
     Ccutlass::gemm::HgemmTransformerA< kLayout_, Iterator_ >
     Ccutlass::gemm::HgemmTransformerA< MatrixLayout::kColumnMajor, Iterator_ >
     Ccutlass::gemm::HgemmTransformerA< MatrixLayout::kRowMajor, Iterator_ >
     Ccutlass::gemm::HgemmTransformerB< kLayout_, Iterator_ >
     Ccutlass::gemm::HgemmTransformerB< MatrixLayout::kColumnMajor, Iterator_ >
     Ccutlass::gemm::HgemmTransformerB< MatrixLayout::kRowMajor, Iterator_ >
     Ccutlass::IdentityDescribes identity elements
     Ccutlass::gemm::IdentityBlockSwizzle
     Ccutlass::IdentityTensorMapFunc< Rank >
     Ccutlass::IdentityTensorMapFunc< Rank_ >
     Ccutlass::gemm::IgemmEpilogueScalar< ScalarD_ >
     Ccutlass::gemm::IgemmEpilogueScalar< int >
     Ccutlass::gemm::IgemmFloatToInt8Converter< kElements_ >
     Ccutlass::gemm::IgemmGlobalLoadTransformer< InputFragment_, OutputScalar_ >
     Ccutlass::gemm::IgemmGlobalLoadTransformer< Fragment< int8_t, kElements_ >, float >
     Ccutlass::gemm::IgemmGlobalStoreTransformer< InputScalar_, OutputFragment_ >
     Ccutlass::gemm::IgemmGlobalStoreTransformer< float, Fragment< int8_t, kElements_ > >
     Ccutlass::gemm::IgemmInt8ToFloatConverter< kElements_ >
     Ccutlass::gemm::IgemmSharedStoreTransformer< InputScalar_, OutputFragment_ >
     Ccutlass::gemm::IgemmSwizzle< GlobalIterator_ >
     Ccutlass::gemm::IgemmTileTraitsHelperA< MatrixLayout::kRowMajor, GemmConfig_, Index_ >
     Ccutlass::gemm::IgemmTileTraitsHelperB< MatrixLayout::kColumnMajor, GemmConfig_, Index_ >
     Ccutlass::gemm::IgemmTraitsHelper< kLayoutA_, kLayoutB_, OutputTile_, ScalarD_, EpilogueFunctor_, ThreadGemmShape_, Index_ >
     Ccutlass::gemm::IgemmTransformerA< kLayout_, Iterator_ >
     Ccutlass::gemm::IgemmTransformerA< MatrixLayout::kColumnMajor, Iterator_ >
     Ccutlass::gemm::IgemmTransformerA< MatrixLayout::kRowMajor, Iterator_ >
     Ccutlass::gemm::IgemmTransformerB< kLayout_, Iterator_ >
     Ccutlass::gemm::IgemmTransformerB< MatrixLayout::kColumnMajor, Iterator_ >
     Ccutlass::gemm::IgemmTransformerB< MatrixLayout::kRowMajor, Iterator_ >
     Ccutlass::int4_t
     ▶Ccutlass::platform::integral_constant< value_t, V >Std::integral_constant
     ▶Ccutlass::platform::integral_constant< bool, V >
     ▶Ccutlass::platform::integral_constant< bool,(is_arithmetic< T >::value||is_void< T >::value||is_same< nullptr_t, remove_cv< T >::type >::value)>
     ▶Ccutlass::platform::integral_constant< bool,(is_base_of_helper< remove_cv< BaseT >::type, remove_cv< DerivedT >::type >::value)||(is_same< remove_cv< BaseT >::type, remove_cv< DerivedT >::type >::value)>
     ▶Ccutlass::platform::integral_constant< bool,(is_fundamental< T >::value||is_pointer< T >::value)>
     ▶Ccutlass::platform::integral_constant< bool,(is_integral< T >::value||is_floating_point< T >::value)>
     ▶Ccutlass::platform::integral_constant< bool,(is_same< float, remove_cv< T >::type >::value||is_same< double, remove_cv< T >::type >::value)>
     ▶Ccutlass::platform::integral_constant< bool,(N &(N - 1))==0 >
     Ccutlass::platform::is_base_of_helper< BaseT, DerivedT >Helper for std::is_base_of
     Ccutlass::PredicateVector< kPredicates_, kPredicatesPerByte_, kPredicateStart_ >::IteratorAn iterator implementing Predicate Iterator Concept enabling sequential read and write access to predicates
     Ccutlass::IteratorAdvanceSpecifies dimension in which post-increment accesses advance
     ▶Ccutlass::KernelLaunchConfigurationStructure containing the basic launch configuration of a CUDA kernel
     Ccutlass::gemm::Launch< Gemm, WithLaunchBounds >Partial specialization for launching the GEMM kernel with or without launch bounds
     Ccutlass::gemm::Launch< Gemm, false >Partial specialization for launching the GEMM kernel with or without launch bounds
     Ccutlass::platform::less< T >Std::less
     ▶Ccutlass::gemm::LinearScaling< Scalar_, FragmentMultiplyAdd_ >Functor to compute linear combination of fragments
     Ccutlass::Load< Scalar_, kAccessSize, Memory_, kFragmentElementType, FragmentElement_, kStride, size >
     Ccutlass::Load< double, 2, Memory_, FragmentElementType::kScalar, double, kStride, 16 >
     Ccutlass::Load< Scalar_, kAccessSize, Memory_, FragmentElementType::kScalar, Scalar_, 1, 2 >Partial specialization for 16b loads
     Ccutlass::Load< Scalar_, kAccessSize, Memory_, FragmentElementType::kScalar, Scalar_, kStride, 16 >
     Ccutlass::Load< Scalar_, kAccessSize, Memory_, FragmentElementType::kScalar, Scalar_, kStride, 4 >
     Ccutlass::Load< Scalar_, kAccessSize, Memory_, FragmentElementType::kScalar, Scalar_, kStride, 8 >
     Ccutlass::Load< Scalar_, kAccessSize, Memory_, FragmentElementType::kWmmaMatrix, FragmentElement_, kStride, size >
     Ccutlass::Load< Vector< bin1_t, 32 >, kAccessSize, Memory_, FragmentElementType::kWmmaMatrix, FragmentElement_, kStride, size >
     Ccutlass::Load< Vector< int4_t, 8 >, kAccessSize, Memory_, FragmentElementType::kWmmaMatrix, FragmentElement_, kStride, size >
     Ccutlass::Load< Vector< uint4_t, 8 >, kAccessSize, Memory_, FragmentElementType::kWmmaMatrix, FragmentElement_, kStride, size >
     Ccutlass::log2_down< N, CurrentVal, Count >
     Ccutlass::log2_down< N, 1, Count >
     Ccutlass::log2_up< N, CurrentVal, Count >
     Ccutlass::log2_up< N, 1, Count >
     Ccutlass::gemm::GemmTraits< GemmConfig_, GlobalLoadStreamA_, GlobalLoadStreamB_, SharedLoadStreamA_, SharedLoadStreamB_, Epilogue_, BlockSwizzle_, Index_, ClearAccumulators_ >::MainLoopSharedStorage
     Ccutlass::MatrixLayoutDefines data layouts of various matrix formats usable by TensorRef and other classes
     Ccutlass::MatrixTransformTransformation applied to matrix operands
     Ccutlass::Max< A, B >
     Ccutlass::MemorySpaceEnum to specify which memory space data resides in
     Ccutlass::Min< A, B >
     Ccutlass::platform::nullptr_tStd::nullptr_t
     Ccutlass::platform::alignment_of< value_t >::pad
     Ccutlass::gemm::LinearScalingDevicePtr< Scalar_, FragmentMultiplyAdd_ >::ParamsThe parameters
     Ccutlass::gemm::GlobalLoadStream< Operand, LoadIterator_, StoreIterator_, Transformer_ >::ParamsThe params
     Ccutlass::gemm::SharedStreamPair< StreamA_, StreamB_ >::ParamsParameters object passed to load iterators
     Ccutlass::ZipTileIterator< First_, Second_ >::ParamsParams object
     Ccutlass::gemm::LinearScaling< Scalar_, FragmentMultiplyAdd_ >::ParamsThe parameters
     Ccutlass::gemm::GlobalLoadStreamPair< StreamA_, StreamB_, kResidueInProlog_ >::ParamsParameters object
     ▶Ccutlass::gemm::GemmGlobalIteratorCd< TileTraits_, Index_ >::ParamsThe params
     Ccutlass::gemm::GemmEpilogueTraits< OutputTile_, Accumulators_, GlobalLoadIteratorC_, GlobalTransformerC_, GlobalTransformerD_, GlobalStoreIteratorD_, SharedStoreIteratorD_, SharedStoreTransformerD_, SharedLoadStreamD_, Iterations_, Delta_, Functor_, Index_ >::ParamsThe params
     ▶Ccutlass::TileIteratorBase< Traits_, Scalar_, Advance_, MemorySpace, Index_, FragmentElement_, FragmentElementType_, Skew_ >::ParamsParameters to the iterator
     Ccutlass::TileLoadStream< Iterator_, Transformer_ >::ParamsParameters object used to construct generic load stream
     Ccutlass::TileStoreStream< Iterator_, Transformer_ >::ParamsParameters used to construct the stream
     Ccutlass::gemm::SharedLoadStream< Iterator_, Transformer_ >::ParamsThe params
     Ccutlass::platform::plus< T >Platform::plus
     Ccutlass::PredicateTileAdapter< PredicateVector_, Iterations_ >Adapter to enable random access to predicates via logical coordinate within a tile
     Ccutlass::TileLoadStream< Iterator_, Transformer_ >::PredicateVectorEmpty predicate vector struct
     Ccutlass::PredicateVector< kPredicates_, kPredicatesPerByte_, kPredicateStart_ >Statically sized array of bits implementing
     Ccutlass::TileStoreStream< Iterator_, Transformer_ >::PredicateVectorEmpty predicate vector struct
     Ccutlass::PredicateVector< Base::Iterations::kW >
     Ccutlass::PredicateVector< ShapeCount< typename Base::Iterations >::kCount >
     Ccutlass::gemm::ProjectOperand< operand, Kstrided >
     Ccutlass::gemm::ProjectOperand< GemmOperand::kA, Kstrided >Project A operand - (0, K, M)
     Ccutlass::gemm::ProjectOperand< GemmOperand::kB, Kstrided >Project B operand - (0, K, N)
     Ccutlass::gemm::ProjectOperand< GemmOperand::kC, true >Project C operand - (0, N, M)
     Ccutlass::gemm::ProjectOperand< GemmOperand::kD, true >Project D operand - (0, N, M)
     Ccutlass::RegularTilePredicateFunctor< Delta_ >Functor computing a predicate given the logical position of an access
     Ccutlass::platform::remove_const< T >Std::remove_const (non-const specialization)
     Ccutlass::platform::remove_const< const T >Std::remove_const (const specialization)
     Ccutlass::platform::remove_cv< T >Std::remove_cv
     Ccutlass::platform::remove_volatile< T >Std::remove_volatile (non-volatile specialization)
     Ccutlass::platform::remove_volatile< volatile T >Std::remove_volatile (volatile specialization)
     Ccutlass::gemm::ReshapeThreads< Tile_, Threads_, bool >
     Ccutlass::gemm::ReshapeThreads< Tile_, Threads_, true >
     Ccutlass::ReshapeTile< Tile_, kAccessSize_, bool >
     Ccutlass::ReshapeTile< Tile_, kAccessSize_, true >
     Ccutlass::MatrixLayout::RowMajorMapping function for row-major matrices
     Ccutlass::MatrixLayout::RowMajorBlockLinear< BlockRows, BlockColumns >
     Ccutlass::gemm::RowMajorBlockSwizzle< groupRows, swDirection >
     Ccutlass::MatrixLayout::RowMajorInterleaved< Interleave >
     Ccutlass::ScalarIO< T >Helper to enable formatted printing of CUTLASS scalar types to an ostream
     Ccutlass::detail::ScalarOrPointer< Scalar_ >
     Ccutlass::detail::ScalarOrPointer< Scalar >
     Ccutlass::Shape< kD_, kH_, kW_, kC_ >A Shape implementing Layout Concept describing the dimensions of a cube
     Ccutlass::ShapeAdd< A_, B_ >
     Ccutlass::ShapeCount< Shape >Compute derived counted of a Layout Concept based class
     Ccutlass::ShapeDiv< A_, B_ >
     Ccutlass::ShapeDivCeiling< A_, B_ >
     Ccutlass::ShapeMax< A_, B_ >
     Ccutlass::ShapeMin< A_, B_ >
     Ccutlass::ShapeMul< A_, B_ >
     Ccutlass::ShapeScale< A_, kScale_ >
     Ccutlass::ShapeStrides< Shape_, elementsPerAccess >
     Ccutlass::ShapeSub< A_, B_ >
     Ccutlass::gemm::SharedLoadStream< Iterator_, Transformer_ >
     Ccutlass::gemm::GemmEpilogueTraits< OutputTile_, Accumulators_, GlobalLoadIteratorC_, GlobalTransformerC_, GlobalTransformerD_, GlobalStoreIteratorD_, SharedStoreIteratorD_, SharedStoreTransformerD_, SharedLoadStreamD_, Iterations_, Delta_, Functor_, Index_ >::SharedStorageThe shared memory to swizzle the data in the epilogue
     Ccutlass::gemm::GlobalLoadStreamPair< StreamA_, StreamB_, kResidueInProlog_ >::SharedStorageDefines a structure containing shared storage for each pair
     Ccutlass::gemm::GemmTraits< GemmConfig_, GlobalLoadStreamA_, GlobalLoadStreamB_, SharedLoadStreamA_, SharedLoadStreamB_, Epilogue_, BlockSwizzle_, Index_, ClearAccumulators_ >::SharedStorageThe storage in shared memory
     Ccutlass::gemm::GlobalLoadStream< Operand, LoadIterator_, StoreIterator_, Transformer_ >::SharedStorage
     Ccutlass::gemm::ClearAccumulators< Scalar_, kLanes_ >::SharedStorageThe shared storage
     Ccutlass::gemm::SharedStreamPair< StreamA_, StreamB_ >Collect the global load streams for multiplicands
     Ccutlass::gemm::SimplifiedGemmTraitsHelper< GemmTileTraitsHelperA_, GemmTileTraitsHelperB_, Index_ >
     Ccutlass::sqrt_est< N >
     Ccutlass::StorageType< alignment >
     Ccutlass::StorageType< 1 >
     Ccutlass::StorageType< 2 >
     Ccutlass::StorageType< 4 >
     Ccutlass::StorageType< kAlignment_ >
     Ccutlass::StorageType< sizeof(Scalar)>
     Ccutlass::Store< Scalar_, kAccessSize, Memory_, kFragmentElementType, FragmentElement_, kStride, size >
     Ccutlass::Store< double, 2, Memory_, FragmentElementType::kScalar, double, kStride, 16 >
     Ccutlass::Store< Scalar_, kAccessSize, Memory_, FragmentElementType::kScalar, Scalar_, 1, 2 >
     Ccutlass::Store< Scalar_, kAccessSize, Memory_, FragmentElementType::kScalar, Scalar_, kStride, 16 >
     Ccutlass::Store< Scalar_, kAccessSize, Memory_, FragmentElementType::kScalar, Scalar_, kStride, 4 >
     Ccutlass::Store< Scalar_, kAccessSize, Memory_, FragmentElementType::kScalar, Scalar_, kStride, 8 >
     Ccutlass::Store< Scalar_, kAccessSize, Memory_, FragmentElementType::kWmmaMatrix, FragmentElement_, kStride, size >
     Ccutlass::gemm::GemmEpilogueTraits< OutputTile_, Accumulators_, GlobalLoadIteratorC_, GlobalTransformerC_, GlobalTransformerD_, GlobalStoreIteratorD_, SharedStoreIteratorD_, SharedStoreTransformerD_, SharedLoadStreamD_, Iterations_, Delta_, Functor_, Index_ >::StreamSharedStorageThe shared memory storage to exchange data
     Ccutlass::TensorRef< Storage_, Rank_, MapFunc_, 1, Index_, LongIndex_ >::StrideVector
     Ccutlass::gemm::swizzleDirection
     ▶Ccutlass::TensorRef< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >
     Ccutlass::TensorRef< AType const, 2 >
     Ccutlass::TensorRef< BType const, 2 >
     Ccutlass::TensorRef< CType const, 2 >
     Ccutlass::TensorRef< DType, 2 >
     Ccutlass::TensorRef< Storage_, Rank_, MapFunc_, 1, Index_, LongIndex_ >Specialization for rank=1 case with no internal StrideVector
     Ccutlass::TensorRefArray< Storage_, Rank_, MapFunc_, StorageRank_, Index_, LongIndex_ >
     Ccutlass::gemm::ThreadMultiplyAdd< ThreadGemmShape_, ThreadsPerWarp_, ScalarA_, ScalarB_, ScalarC_, kLayout_ >Template performing matrix multiply-add operation within a thread
     Ccutlass::gemm::ThreadMultiplyAdd< ThreadGemmShape_, ThreadsPerWarp_, half, half, float >Template performing matrix multiply-add operation within a thread
     Ccutlass::gemm::ThreadMultiplyAdd< ThreadGemmShape_, ThreadsPerWarp_, half, half, half >Template performing matrix multiply-add operation within a thread
     Ccutlass::gemm::ThreadMultiplyAdd< ThreadGemmShape_, ThreadsPerWarp_, int8_t, int8_t, int >Template performing matrix multiply-add operation within a thread
     Ccutlass::gemm::GemmSharedStoreTileAbTraits< Scalar_, Tile_, Threads_, kScalarsPerSts_ >::ThreadOffset
     Ccutlass::gemm::WmmaGemmGlobalIteratorCdTraits< Scalar_, Tile_, Threads_, kAccessSize_ >::ThreadOffsetComputes the thread offset in (H, W) based on thread ID
     Ccutlass::gemm::GemmGlobalTileCdTraits< Scalar_, Tile_, Threads_, kStrideH_, kAccessSize_ >::ThreadOffsetComputes the thread offset in (H, W) based on thread ID
     Ccutlass::gemm::GemmSharedLoadTileATraits< Scalar_, OutputTile_, Warps_, ThreadsPerWarp_, InstructionShape_, kStages_, kScalarsPerLds_, kSkew_ >::ThreadOffsetComputes the thread offset in (H, W) based on thread ID
     Ccutlass::gemm::GemmSharedStoreWithSkewTileAbTraits< Scalar_, Tile_, Threads_, kScalarsPerSts_, kSkew_ >::ThreadOffset
     Ccutlass::gemm::IgemmGlobalTileTraits< kOperand_, kLayout_, Scalar_, Tile_, Threads_, kAccessSize_ >::ThreadOffsetComputes the thread offset in (H, W) based on thread ID
     Ccutlass::gemm::GemmSharedLoadTileBTraits< Scalar_, OutputTile_, Warps_, ThreadsPerWarp_, InstructionShape_, kStages_, kScalarsPerLds_, kSkew_ >::ThreadOffsetComputes the thread offset in (H, W) based on thread ID
     Ccutlass::gemm::GemmGlobalTileTraits< kOperand_, kLayout_, Scalar_, Tile_, Threads_, kAccessSize_ >::ThreadOffsetComputes the thread offset in (H, W) based on thread ID
     Ccutlass::gemm::GemmSharedLoadTileDTraits< Scalar_, OutputTile_, Warps_, ThreadsPerWarp_, kTileH_, kScalarsPerLds_, kSkew_ >::ThreadOffsetComputes the thread offset in (H, W) based on thread ID
     Ccutlass::TileTraitsWarpRake< Tile_, Threads >::ThreadOffsetComputes the thread offset in (H, W) based on thread ID
     Ccutlass::gemm::GemmSharedStoreTileDTraits< Scalar_, OutputTile_, Warps_, ThreadsPerWarp_, kScalarsPerSts_, kSkew_ >::ThreadOffsetComputes the thread offset in (H, W) based on thread ID
     Ccutlass::gemm::HgemmCrosswiseGlobalTileTraits< kOperand_, kLayout_, Scalar_, Tile_, Threads_, kAccessSize_ >::ThreadOffsetComputes the thread offset in (H, W) based on thread ID
     Ccutlass::TileAllocation< Scalar_, Shape_ >Class for storing a tile in memory and accessing it through a tensor ref
     Ccutlass::TiledThreadOffset< ThreadShape >Basic thread offset function computed from a thread shape
     ▶Ccutlass::TileIteratorBase< Traits_, Scalar_, Advance_, MemorySpace, Index_, FragmentElement_, FragmentElementType_, Skew_ >Iterator for accessing a stripmined tile in memory
     ▶Ccutlass::TileIteratorBase< TileTraits_, TileTraits_::Scalar, Advance_, MemorySpace, Index_, TileTraits_::Scalar, FragmentElementType::kScalar, Shape< 0, 0, 0, 0 > >
     ▶Ccutlass::TileIteratorBase< TileTraits_, TileTraits_::Scalar, IteratorAdvance::kH, MemorySpace::kGlobal, Index_ >
     ▶Ccutlass::TileLoadStream< Iterator_, Transformer_ >Generic stream for loading and transforming fragments
     ▶Ccutlass::TileStoreStream< Iterator_, Transformer_ >Generic stream for transforming and storing fragments
     Ccutlass::TileTraits< Tile_, Delta_, Iterations_, ThreadOffset_, AccessSize >A template defining Tile Traits Concept
     Ccutlass::TileTraitsContiguousMajor< Tile_, Threads >
     Ccutlass::TileTraitsStandard< Tile_, Threads >Chooses 'best' shape to enable warp raking along contiguous dimension if possible
     Ccutlass::TileTraitsStrideMajor< Tile_, Threads >
     Ccutlass::TileTraitsWarpRake< Tile_, Threads >Tiling in which warps rake across the contiguous dimension
     Ccutlass::PredicateVector< kPredicates_, kPredicatesPerByte_, kPredicateStart_ >::TrivialIteratorIterator that always returns true
     Ccutlass::TrivialPredicateTileAdapterAlways returns true predicate
     Ccutlass::uint4_t
     Ccutlass::platform::unique_ptr< T, Deleter >Std::unique_ptr
     Ccutlass::Vector< Scalar_, kLanes_ >
     Ccutlass::Vector< bin1_t, kLanes_ >Vector definition for 1-bit binary datatype
     Ccutlass::Vector< half, 1 >
     Ccutlass::Vector< half, kLanes_ >
     Ccutlass::Vector< int4_t, kLanes_ >Vector definition for 4-bit signed integer datatype
     Ccutlass::Vector< uint4_t, kLanes_ >Vector definition for 4-bit unsigned integer datatype
     Ccutlass::Vectorize< Element_, kLanes_ >
     Ccutlass::Vectorize< Vector< bin1_t, 32 >, kLanes_ >
     Ccutlass::Vectorize< Vector< int4_t, 8 >, kLanes_ >
     Ccutlass::Vectorize< Vector< uint4_t, 8 >, kLanes_ >
     Ccutlass::VectorTraits< T >Traits describing properties of vectors and scalar-as-vectors
     Ccutlass::VectorTraits< Vector< T, Lanes > >Partial specialization for actual cutlass::Vector
     Ccutlass::VectorTraits< Vector< T, Lanes > const >Partial specialization for actual cutlass::Vector
     Ccutlass::ZipConvert< First_, Second_ >Zips two convert operations
     Ccutlass::ZipFragment< First_, Second_ >A template defining Fragment Concept
     Ccutlass::ZipTensorRef< First_, Second_ >
     Ccutlass::ZipTileAllocation< First_, Second_ >Manages a pair of tile allocations as if they are one allocation
     Ccutlass::ZipTileIterator< First_, Second_ >Constructs an iterator from a pair of iterators
    -Generated on Thu May 3 2018 16:36:15 for Cutlass by   +Generated on Wed Sep 19 2018 13:58:53 for Cutlass by   doxygen 1.8.14
    diff --git a/docs/igemm__epilogue_8h.html b/docs/igemm__epilogue_8h.html index 9b5e5ccf..f7332de2 100644 --- a/docs/igemm__epilogue_8h.html +++ b/docs/igemm__epilogue_8h.html @@ -82,13 +82,13 @@ $(function() {

    Defines the epilogue phase of the GEMM computation for IGEMM, supporting integer and floating-point output matrix formats. More...

    -
    #include <cutlass/convert.h>
    -#include <cutlass/fragment.h>
    -#include <cutlass/gemm/gemm_global_stream.h>
    -#include <cutlass/gemm/gemm_shared_stream.h>
    -#include <cutlass/gemm/igemm_global_tile.h>
    -#include <cutlass/reshape_tile.h>
    -#include <cutlass/tile_iterator.h>
    +
    #include "cutlass/convert.h"
    +#include "cutlass/fragment.h"
    +#include "cutlass/gemm/gemm_global_stream.h"
    +#include "cutlass/gemm/gemm_shared_stream.h"
    +#include "cutlass/gemm/igemm_global_tile.h"
    +#include "cutlass/reshape_tile.h"
    +#include "cutlass/tile_iterator.h"

    Go to the source code of this file.

    @@ -127,7 +127,7 @@ Namespaces
    -Generated on Thu May 3 2018 16:36:12 for Cutlass by   +Generated on Wed Sep 19 2018 13:58:50 for Cutlass by   doxygen 1.8.14
    diff --git a/docs/igemm__epilogue_8h_source.html b/docs/igemm__epilogue_8h_source.html index bfef820a..43f9f158 100644 --- a/docs/igemm__epilogue_8h_source.html +++ b/docs/igemm__epilogue_8h_source.html @@ -76,67 +76,66 @@ $(function() {
    igemm_epilogue.h
    -Go to the documentation of this file.
    1 /***************************************************************************************************
    2  * Copyright (c) 2017-2018, NVIDIA CORPORATION. All rights reserved.
    3  *
    4  * Redistribution and use in source and binary forms, with or without modification, are permitted
    5  * provided that the following conditions are met:
    6  * * Redistributions of source code must retain the above copyright notice, this list of
    7  * conditions and the following disclaimer.
    8  * * Redistributions in binary form must reproduce the above copyright notice, this list of
    9  * conditions and the following disclaimer in the documentation and/or other materials
    10  * provided with the distribution.
    11  * * Neither the name of the NVIDIA CORPORATION nor the names of its contributors may be used
    12  * to endorse or promote products derived from this software without specific prior written
    13  * permission.
    14  *
    15  * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND ANY EXPRESS OR
    16  * IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND
    17  * FITNESS FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE
    18  * FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING,
    19  * BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS;
    20  * OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT,
    21  * STRICT LIABILITY, OR TOR (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
    22  * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
    23  *
    24  **************************************************************************************************/
    29 #pragma once
    30 
    31 #include <cutlass/convert.h>
    32 #include <cutlass/fragment.h>
    33 #include <cutlass/gemm/gemm_global_stream.h>
    34 #include <cutlass/gemm/gemm_shared_stream.h>
    35 #include <cutlass/gemm/igemm_global_tile.h>
    36 #include <cutlass/reshape_tile.h>
    37 #include <cutlass/tile_iterator.h>
    38 
    39 namespace cutlass {
    40 namespace gemm {
    41 
    43 
    44 template <int kElements_>
    45 struct IgemmFloatToInt8Converter {
    47  typedef Fragment<float, kElements_> InputFragment;
    49  typedef Fragment<int8_t, kElements_> OutputFragment;
    50 
    51  // We are packing 4 floats into int32 registers so we need kElements to be multiple of 4.
    52  static_assert(kElements_ % 4 == 0, "kElements must be multiple of 4");
    53 
    55  CUTLASS_DEVICE IgemmFloatToInt8Converter() {}
    56 
    58  CUTLASS_DEVICE void transform(InputFragment const& src, OutputFragment& dst) {
    59  transform(src, 0, dst);
    60  }
    61 
    63  template <typename Fragment_>
    64  CUTLASS_DEVICE void transform(Fragment_ const& src, int offset, OutputFragment& dst) {
    65  // The inputs.
    66  float4 const* src_f4 = reinterpret_cast<float4 const*>(&src[0]);
    67  // The outputs.
    68  int* dst_int = reinterpret_cast<int*>(&dst[0]);
    69 
    70  // Iterate over the floats and pack them together to produce ints.
    71  for (int i = 0; i < kElements_ / 4; ++i) {
    72  // Read the float4.
    73  float4 f4 = src_f4[i];
    74 
    75  // Clamp the 4 elements of the floats to the [-128, +127] range.
    76  float x = fmaxf(-128.f, fminf(127.f, f4.x));
    77  float y = fmaxf(-128.f, fminf(127.f, f4.y));
    78  float z = fmaxf(-128.f, fminf(127.f, f4.z));
    79  float w = fmaxf(-128.f, fminf(127.f, f4.w));
    80 
    81  // Convert to integers.
    82  int ix = (int)x;
    83  int iy = (int)y;
    84  int iz = (int)z;
    85  int iw = (int)w;
    86 
    87  // Extract the lower bytes to build an int32 with 4 int8.
    88  asm volatile("prmt.b32 %0, %0, %1, 0x1140;" : "+r"(ix) : "r"(iy));
    89  asm volatile("prmt.b32 %0, %0, %1, 0x1140;" : "+r"(iz) : "r"(iw));
    90  asm volatile("prmt.b32 %0, %0, %1, 0x5410;" : "+r"(ix) : "r"(iz));
    91 
    92  // Store the int.
    93  dst_int[i] = ix;
    94  }
    95  }
    96 };
    97 
    99 
    100 template <typename InputScalar_, typename OutputFragment_>
    101 struct IgemmGlobalStoreTransformer {
    102  typedef Convert<Fragment<InputScalar_, OutputFragment_::kElements>, OutputFragment_> Transformer;
    103 };
    104 
    105 template <int kElements_>
    106 struct IgemmGlobalStoreTransformer<float, Fragment<int8_t, kElements_> > {
    107  typedef IgemmFloatToInt8Converter<kElements_> Transformer;
    108 };
    109 
    111 
    112 template <int kElements_>
    113 struct IgemmInt8ToFloatConverter {
    115  typedef Fragment<int8_t, kElements_> InputFragment;
    117  typedef Fragment<float, kElements_> OutputFragment;
    118 
    119  // We are unpacking 4 int8s from int32.
    120  static_assert(kElements_ % 4 == 0, "kElements must be multiple of 4");
    121 
    123  CUTLASS_DEVICE IgemmInt8ToFloatConverter() {}
    124 
    126  CUTLASS_DEVICE void transform(InputFragment const& src, OutputFragment& dst) {
    127  transform(src, 0, dst);
    128  }
    129 
    131  template <typename Fragment_>
    132  CUTLASS_DEVICE void transform(Fragment_ const& src, int offset, OutputFragment& dst) {
    133  // The inputs.
    134  int const* src_int = reinterpret_cast<int const*>(&src[0]);
    135  // The outputs.
    136  float4* dst_f4 = reinterpret_cast<float4*>(&dst[0]);
    137 
    138  // Iterate over the int8 and unpack them together to produce floats.
    139  for (int i = 0; i < kElements_ / 4; ++i) {
    140  // Read the int.
    141  int ix, iy, iz, iw = src_int[i];
    142 
    143  // Extract the 4 bytes.
    144  asm volatile("prmt.b32 %0, 0x0, %1, 0x4440;" : "=r"(ix) : "r"(iw));
    145  asm volatile("prmt.b32 %0, 0x0, %1, 0x4441;" : "=r"(iy) : "r"(iw));
    146  asm volatile("prmt.b32 %0, 0x0, %1, 0x4442;" : "=r"(iz) : "r"(iw));
    147  asm volatile("prmt.b32 %0, 0x0, %1, 0x4443;" : "=r"(iw) : "r"(iw));
    148 
    149  // The floats.
    150  float fx, fy, fz, fw;
    151 
    152  // Convert to floats (make sure we generate I2F.F32.S8).
    153  asm volatile("cvt.rn.f32.s8 %0, %1;" : "=f"(fx) : "r"(ix));
    154  asm volatile("cvt.rn.f32.s8 %0, %1;" : "=f"(fy) : "r"(iy));
    155  asm volatile("cvt.rn.f32.s8 %0, %1;" : "=f"(fz) : "r"(iz));
    156  asm volatile("cvt.rn.f32.s8 %0, %1;" : "=f"(fw) : "r"(iw));
    157 
    158  // Store the float4.
    159  dst_f4[i] = make_float4(fx, fy, fz, fw);
    160  }
    161  }
    162 };
    163 
    165 
    166 template <typename InputFragment_, typename OutputScalar_>
    167 struct IgemmGlobalLoadTransformer {
    168  typedef Convert<InputFragment_, Fragment<OutputScalar_, InputFragment_::kElements> > Transformer;
    169 };
    170 
    171 template <int kElements_>
    172 struct IgemmGlobalLoadTransformer<Fragment<int8_t, kElements_>, float> {
    173  typedef IgemmInt8ToFloatConverter<kElements_> Transformer;
    174 };
    175 
    177 
    178 template <typename InputScalar_, typename OutputFragment_>
    179 struct IgemmSharedStoreTransformer {
    180  typedef Convert<Fragment<InputScalar_, OutputFragment_::kElements>, OutputFragment_> Transformer;
    181 };
    182 
    184 
    185 template <typename IgemmConfig_, typename EpilogueFunctor_, typename Index_>
    186 struct IgemmEpilogueTraitsHelper
    187  : public GemmEpilogueTraitsHelper<IgemmConfig_, EpilogueFunctor_, Index_> {
    189  typedef GemmEpilogueTraitsHelper<IgemmConfig_, EpilogueFunctor_, Index_> Base;
    191  typedef IgemmConfig_ IgemmConfig;
    192 
    194  typedef typename Base::Scalar Scalar;
    196  typedef typename Base::Iterations Iterations;
    198  typedef typename Base::Delta Delta;
    199 
    201  typedef typename Base::GlobalLoadTileTraits GlobalLoadTileTraits;
    203  typedef GemmGlobalIteratorCd<GlobalLoadTileTraits> GlobalLoadIteratorC;
    205  typedef typename GlobalLoadIteratorC::Fragment GlobalFragmentC;
    207  typedef
    208  typename IgemmGlobalLoadTransformer<GlobalFragmentC, Scalar>::Transformer GlobalTransformerC;
    209 
    211  typedef typename Base::GlobalStoreTileTraits GlobalStoreTileTraits;
    213  typedef GemmGlobalIteratorCd<GlobalStoreTileTraits> GlobalStoreIteratorD;
    215  typedef typename GlobalStoreIteratorD::Fragment GlobalFragmentD;
    217  typedef
    218  typename IgemmGlobalStoreTransformer<Scalar, GlobalFragmentD>::Transformer GlobalTransformerD;
    219 
    221  typedef typename Base::SharedStoreTileTraits SharedStoreTileTraits;
    223  typedef TileStoreIterator<SharedStoreTileTraits,
    224  typename SharedStoreTileTraits::Scalar,
    225  IteratorAdvance::kH,
    226  MemorySpace::kGlobal>
    227  SharedStoreIteratorD;
    229  typedef typename SharedStoreIteratorD::Fragment SharedStoreFragmentD;
    231  typedef typename IgemmSharedStoreTransformer<typename IgemmConfig::Accumulators::Element,
    232  SharedStoreFragmentD>::Transformer
    233  SharedStoreTransformerD;
    235  typedef typename Base::SharedLoadTileTraits SharedLoadTileTraits;
    237  typedef TileLoadIterator<SharedLoadTileTraits,
    238  typename SharedLoadTileTraits::Scalar,
    239  IteratorAdvance::kH,
    240  MemorySpace::kShared>
    241  SharedLoadIteratorD;
    242 };
    243 
    245 
    246 template <
    248  typename IgemmConfig_,
    250  typename EpilogueFunctor_,
    252  typename Index_ = int,
    254  typename Helper_ = IgemmEpilogueTraitsHelper<IgemmConfig_, EpilogueFunctor_, Index_> >
    255 struct IgemmEpilogueTraits : public GemmEpilogueTraits<
    256  // The output tile.
    257  typename IgemmConfig_::OutputTile,
    258  // The accumulators.
    259  typename IgemmConfig_::Accumulators,
    260  // The global iterator for C.
    261  typename Helper_::GlobalLoadIteratorC,
    262  // The transformer for C.
    263  typename Helper_::GlobalTransformerC,
    264  // The transformer for D.
    265  typename Helper_::GlobalTransformerD,
    266  // The global iterator for D.
    267  typename Helper_::GlobalStoreIteratorD,
    268  // The iterator to store D to shared memory.
    269  typename Helper_::SharedStoreIteratorD,
    270  // The shared store transformer for D.
    271  typename Helper_::SharedStoreTransformerD,
    272  // The iterator to load D from shared memory.
    273  typename Helper_::SharedLoadIteratorD,
    274  // The iterations.
    275  typename Helper_::Iterations,
    276  // The strides between iterations.
    277  typename Helper_::Delta,
    278  // The functor to be used in the epilogue.
    279  EpilogueFunctor_,
    280  // The index.
    281  Index_> {
    283  static bool const kInt8Output =
    284  platform::is_same<typename IgemmConfig_::ScalarC, int8_t>::value != 0;
    285 };
    286 
    288 
    289 template <typename GemmEpilogueTraits_, bool = GemmEpilogueTraits_::kInt8Output>
    290 struct IgemmEpilogue : public GemmEpilogue<GemmEpilogueTraits_> {
    292  typedef GemmEpilogue<GemmEpilogueTraits_> Base;
    293 
    295  CUTLASS_DEVICE IgemmEpilogue(typename Base::Params const& params_,
    296  typename Base::SharedStorage& shared_storage_,
    297  typename Base::Index m_,
    298  typename Base::Index n_)
    299  : Base(params_, shared_storage_, m_, n_) {}
    300 };
    301 
    303 
    304 template <typename GemmEpilogueTraits_>
    305 struct IgemmEpilogue<GemmEpilogueTraits_, true> : public GemmEpilogue<GemmEpilogueTraits_> {
    307  typedef GemmEpilogue<GemmEpilogueTraits_> Base;
    308 
    310  CUTLASS_DEVICE IgemmEpilogue(typename Base::Params const& params_,
    311  typename Base::SharedStorage& shared_storage_,
    312  typename Base::Index m_,
    313  typename Base::Index n_)
    314  : Base(params_, shared_storage_, m_, n_) {}
    315 };
    316 
    318 
    319 } // namespace gemm
    320 } // namespace cutlass
    cutlass::gemm::GemmGlobalTileCdTraits
    Definition: gemm_global_tile.h:116
    +Go to the documentation of this file.
    1 /***************************************************************************************************
    2  * Copyright (c) 2017-2018, NVIDIA CORPORATION. All rights reserved.
    3  *
    4  * Redistribution and use in source and binary forms, with or without modification, are permitted
    5  * provided that the following conditions are met:
    6  * * Redistributions of source code must retain the above copyright notice, this list of
    7  * conditions and the following disclaimer.
    8  * * Redistributions in binary form must reproduce the above copyright notice, this list of
    9  * conditions and the following disclaimer in the documentation and/or other materials
    10  * provided with the distribution.
    11  * * Neither the name of the NVIDIA CORPORATION nor the names of its contributors may be used
    12  * to endorse or promote products derived from this software without specific prior written
    13  * permission.
    14  *
    15  * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND ANY EXPRESS OR
    16  * IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND
    17  * FITNESS FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE
    18  * FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING,
    19  * BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS;
    20  * OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT,
    21  * STRICT LIABILITY, OR TOR (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
    22  * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
    23  *
    24  **************************************************************************************************/
    29 #pragma once
    30 
    31 #include "cutlass/convert.h"
    32 #include "cutlass/fragment.h"
    33 #include "cutlass/gemm/gemm_global_stream.h"
    34 #include "cutlass/gemm/gemm_shared_stream.h"
    35 #include "cutlass/gemm/igemm_global_tile.h"
    36 #include "cutlass/reshape_tile.h"
    37 #include "cutlass/tile_iterator.h"
    38 
    39 namespace cutlass {
    40 namespace gemm {
    41 
    43 
    44 template <int kElements_>
    45 struct IgemmFloatToInt8Converter {
    47  typedef Fragment<float, kElements_> InputFragment;
    49  typedef Fragment<int8_t, kElements_> OutputFragment;
    50 
    51  // We are packing 4 floats into int32 registers so we need kElements to be multiple of 4.
    52  static_assert(kElements_ % 4 == 0, "kElements must be multiple of 4");
    53 
    55  CUTLASS_DEVICE IgemmFloatToInt8Converter() {}
    56 
    58  CUTLASS_DEVICE void transform(InputFragment const& src, OutputFragment& dst) {
    59  transform(src, 0, dst);
    60  }
    61 
    63  template <typename Fragment_>
    64  CUTLASS_DEVICE void transform(Fragment_ const& src, int offset, OutputFragment& dst) {
    65  // The inputs.
    66  float4 const* src_f4 = reinterpret_cast<float4 const*>(&src[0]);
    67  // The outputs.
    68  int* dst_int = reinterpret_cast<int*>(&dst[0]);
    69 
    70  // Iterate over the floats and pack them together to produce ints.
    71  for (int i = 0; i < kElements_ / 4; ++i) {
    72  // Read the float4.
    73  float4 f4 = src_f4[i];
    74 
    75  // Clamp the 4 elements of the floats to the [-128, +127] range.
    76  float x = fmaxf(-128.f, fminf(127.f, f4.x));
    77  float y = fmaxf(-128.f, fminf(127.f, f4.y));
    78  float z = fmaxf(-128.f, fminf(127.f, f4.z));
    79  float w = fmaxf(-128.f, fminf(127.f, f4.w));
    80 
    81  // Convert to integers.
    82  int ix = (int)x;
    83  int iy = (int)y;
    84  int iz = (int)z;
    85  int iw = (int)w;
    86 
    87  // Extract the lower bytes to build an int32 with 4 int8.
    88  asm volatile("prmt.b32 %0, %0, %1, 0x1140;" : "+r"(ix) : "r"(iy));
    89  asm volatile("prmt.b32 %0, %0, %1, 0x1140;" : "+r"(iz) : "r"(iw));
    90  asm volatile("prmt.b32 %0, %0, %1, 0x5410;" : "+r"(ix) : "r"(iz));
    91 
    92  // Store the int.
    93  dst_int[i] = ix;
    94  }
    95  }
    96 };
    97 
    99 
    100 template <typename InputScalar_, typename OutputFragment_>
    101 struct IgemmGlobalStoreTransformer {
    102  typedef Convert<Fragment<InputScalar_, OutputFragment_::kElements>, OutputFragment_> Transformer;
    103 };
    104 
    105 template <int kElements_>
    106 struct IgemmGlobalStoreTransformer<float, Fragment<int8_t, kElements_> > {
    107  typedef IgemmFloatToInt8Converter<kElements_> Transformer;
    108 };
    109 
    111 
    112 template <int kElements_>
    113 struct IgemmInt8ToFloatConverter {
    115  typedef Fragment<int8_t, kElements_> InputFragment;
    117  typedef Fragment<float, kElements_> OutputFragment;
    118 
    119  // We are unpacking 4 int8s from int32.
    120  static_assert(kElements_ % 4 == 0, "kElements must be multiple of 4");
    121 
    123  CUTLASS_DEVICE IgemmInt8ToFloatConverter() {}
    124 
    126  CUTLASS_DEVICE void transform(InputFragment const& src, OutputFragment& dst) {
    127  transform(src, 0, dst);
    128  }
    129 
    131  template <typename Fragment_>
    132  CUTLASS_DEVICE void transform(Fragment_ const& src, int offset, OutputFragment& dst) {
    133  // The inputs.
    134  int const* src_int = reinterpret_cast<int const*>(&src[0]);
    135  // The outputs.
    136  float4* dst_f4 = reinterpret_cast<float4*>(&dst[0]);
    137 
    138  // Iterate over the int8 and unpack them together to produce floats.
    139  for (int i = 0; i < kElements_ / 4; ++i) {
    140  // Read the int.
    141  int ix, iy, iz, iw = src_int[i];
    142 
    143  // Extract the 4 bytes.
    144  asm volatile("prmt.b32 %0, 0x0, %1, 0x4440;" : "=r"(ix) : "r"(iw));
    145  asm volatile("prmt.b32 %0, 0x0, %1, 0x4441;" : "=r"(iy) : "r"(iw));
    146  asm volatile("prmt.b32 %0, 0x0, %1, 0x4442;" : "=r"(iz) : "r"(iw));
    147  asm volatile("prmt.b32 %0, 0x0, %1, 0x4443;" : "=r"(iw) : "r"(iw));
    148 
    149  // The floats.
    150  float fx, fy, fz, fw;
    151 
    152  // Convert to floats (make sure we generate I2F.F32.S8).
    153  asm volatile("cvt.rn.f32.s8 %0, %1;" : "=f"(fx) : "r"(ix));
    154  asm volatile("cvt.rn.f32.s8 %0, %1;" : "=f"(fy) : "r"(iy));
    155  asm volatile("cvt.rn.f32.s8 %0, %1;" : "=f"(fz) : "r"(iz));
    156  asm volatile("cvt.rn.f32.s8 %0, %1;" : "=f"(fw) : "r"(iw));
    157 
    158  // Store the float4.
    159  dst_f4[i] = make_float4(fx, fy, fz, fw);
    160  }
    161  }
    162 };
    163 
    165 
    166 template <typename InputFragment_, typename OutputScalar_>
    167 struct IgemmGlobalLoadTransformer {
    168  typedef Convert<InputFragment_, Fragment<OutputScalar_, InputFragment_::kElements> > Transformer;
    169 };
    170 
    171 template <int kElements_>
    172 struct IgemmGlobalLoadTransformer<Fragment<int8_t, kElements_>, float> {
    173  typedef IgemmInt8ToFloatConverter<kElements_> Transformer;
    174 };
    175 
    177 
    178 template <typename InputScalar_, typename OutputFragment_>
    179 struct IgemmSharedStoreTransformer {
    180  typedef Convert<Fragment<InputScalar_, OutputFragment_::kElements>, OutputFragment_> Transformer;
    181 };
    182 
    184 
    185 template <typename IgemmConfig_, typename EpilogueFunctor_, typename Index_>
    186 struct IgemmEpilogueTraitsHelper
    187  : public GemmEpilogueTraitsHelper<IgemmConfig_, EpilogueFunctor_, Index_> {
    189  typedef GemmEpilogueTraitsHelper<IgemmConfig_, EpilogueFunctor_, Index_> Base;
    191  typedef IgemmConfig_ IgemmConfig;
    192 
    194  typedef typename Base::Scalar Scalar;
    196  typedef typename Base::Iterations Iterations;
    198  typedef typename Base::Delta Delta;
    199 
    201  typedef typename Base::GlobalLoadTileTraits GlobalLoadTileTraits;
    203  typedef GemmGlobalIteratorCd<GlobalLoadTileTraits> GlobalLoadIteratorC;
    205  typedef typename GlobalLoadIteratorC::Fragment GlobalFragmentC;
    207  typedef
    208  typename IgemmGlobalLoadTransformer<GlobalFragmentC, Scalar>::Transformer GlobalTransformerC;
    209 
    211  typedef typename Base::GlobalStoreTileTraits GlobalStoreTileTraits;
    213  typedef GemmGlobalIteratorCd<GlobalStoreTileTraits> GlobalStoreIteratorD;
    215  typedef typename GlobalStoreIteratorD::Fragment GlobalFragmentD;
    217  typedef
    218  typename IgemmGlobalStoreTransformer<Scalar, GlobalFragmentD>::Transformer GlobalTransformerD;
    219 
    221  typedef typename Base::SharedStoreTileTraits SharedStoreTileTraits;
    223  typedef TileStoreIterator<SharedStoreTileTraits,
    224  typename SharedStoreTileTraits::Scalar,
    225  IteratorAdvance::kH,
    226  MemorySpace::kGlobal>
    227  SharedStoreIteratorD;
    229  typedef typename SharedStoreIteratorD::Fragment SharedStoreFragmentD;
    231  typedef typename IgemmSharedStoreTransformer<typename IgemmConfig::Accumulators::Element,
    232  SharedStoreFragmentD>::Transformer
    233  SharedStoreTransformerD;
    235  typedef typename Base::SharedLoadTileTraits SharedLoadTileTraits;
    237  typedef TileLoadIterator<SharedLoadTileTraits,
    238  typename SharedLoadTileTraits::Scalar,
    239  IteratorAdvance::kH,
    240  MemorySpace::kShared>
    241  SharedLoadIteratorD;
    242 };
    243 
    245 
    246 template <
    248  typename IgemmConfig_,
    250  typename EpilogueFunctor_,
    252  typename Index_ = int,
    254  typename Helper_ = IgemmEpilogueTraitsHelper<IgemmConfig_, EpilogueFunctor_, Index_> >
    255 struct IgemmEpilogueTraits : public GemmEpilogueTraits<
    256  // The output tile.
    257  typename IgemmConfig_::OutputTile,
    258  // The accumulators.
    259  typename IgemmConfig_::Accumulators,
    260  // The global iterator for C.
    261  typename Helper_::GlobalLoadIteratorC,
    262  // The transformer for C.
    263  typename Helper_::GlobalTransformerC,
    264  // The transformer for D.
    265  typename Helper_::GlobalTransformerD,
    266  // The global iterator for D.
    267  typename Helper_::GlobalStoreIteratorD,
    268  // The iterator to store D to shared memory.
    269  typename Helper_::SharedStoreIteratorD,
    270  // The shared store transformer for D.
    271  typename Helper_::SharedStoreTransformerD,
    272  // The stream to load D from shared memory.
    273  typename Helper_::SharedLoadStreamD,
    274  // The iterations.
    275  typename Helper_::Iterations,
    276  // The strides between iterations.
    277  typename Helper_::Delta,
    278  // The functor to be used in the epilogue.
    279  EpilogueFunctor_,
    280  // The index.
    281  Index_> {
    283  static bool const kInt8Output =
    284  platform::is_same<typename IgemmConfig_::ScalarC, int8_t>::value != 0;
    285 };
    286 
    288 
    289 template <typename GemmEpilogueTraits_, bool = GemmEpilogueTraits_::kInt8Output>
    290 struct IgemmEpilogue : public GemmEpilogue<GemmEpilogueTraits_> {
    292  typedef GemmEpilogue<GemmEpilogueTraits_> Base;
    293 
    295  CUTLASS_DEVICE IgemmEpilogue(typename Base::Params const& params_,
    296  typename Base::SharedStorage& shared_storage_,
    297  Coord<3> const& _problem_size)
    298  : Base(params_, shared_storage_, _problem_size) {}
    299 };
    300 
    302 
    303 template <typename GemmEpilogueTraits_>
    304 struct IgemmEpilogue<GemmEpilogueTraits_, true> : public GemmEpilogue<GemmEpilogueTraits_> {
    306  typedef GemmEpilogue<GemmEpilogueTraits_> Base;
    307 
    309  CUTLASS_DEVICE IgemmEpilogue(typename Base::Params const& params_,
    310  typename Base::SharedStorage& shared_storage_,
    311  Coord<3> const& _problem_size)
    312  : Base(params_, shared_storage_, _problem_size) {}
    313 };
    314 
    316 
    317 } // namespace gemm
    318 } // namespace cutlass
    cutlass::gemm::GemmGlobalTileCdTraits
    Definition: gemm_global_tile.h:120
    cutlass::gemm::IgemmEpilogueTraits
    Definition: igemm_epilogue.h:255
    -
    cutlass::MemorySpace::kShared
    Definition: load_store.h:42
    +
    cutlass::MemorySpace::kShared
    Definition: load_store.h:41
    cutlass::gemm::IgemmEpilogueTraitsHelper::Delta
    Base::Delta Delta
    The iterations strides.
    Definition: igemm_epilogue.h:198
    -
    cutlass::TileStoreIterator::Fragment
    Base::Fragment Fragment
    Fragment definition.
    Definition: tile_iterator.h:682
    cutlass::gemm::IgemmEpilogueTraitsHelper::SharedStoreTileTraits
    Base::SharedStoreTileTraits SharedStoreTileTraits
    The traits class for the shared iterator to store D to shared memory.
    Definition: igemm_epilogue.h:221
    cutlass::gemm::IgemmEpilogueTraitsHelper::GlobalTransformerD
    IgemmGlobalStoreTransformer< Scalar, GlobalFragmentD >::Transformer GlobalTransformerD
    The transformer from accumulators to shared memory fragments.
    Definition: igemm_epilogue.h:218
    cutlass
    Definition: convert.h:33
    cutlass::gemm::IgemmEpilogueTraitsHelper::SharedLoadTileTraits
    Base::SharedLoadTileTraits SharedLoadTileTraits
    The traits class for the shared iterator to load D from shared memory.
    Definition: igemm_epilogue.h:235
    cutlass::gemm::IgemmEpilogueTraitsHelper::SharedLoadIteratorD
    TileLoadIterator< SharedLoadTileTraits, typename SharedLoadTileTraits::Scalar, IteratorAdvance::kH, MemorySpace::kShared > SharedLoadIteratorD
    The shared iterator to load D from shared memory.
    Definition: igemm_epilogue.h:241
    -
    cutlass::gemm::GemmEpilogueTraitsHelper
    Definition: gemm_epilogue_traits.h:171
    +
    cutlass::gemm::GemmEpilogueTraitsHelper
    Definition: gemm_epilogue_traits.h:186
    cutlass::gemm::IgemmEpilogue::Base
    GemmEpilogue< GemmEpilogueTraits_ > Base
    The base class.
    Definition: igemm_epilogue.h:292
    -
    cutlass::gemm::GemmEpilogue::Params
    Traits::Params Params
    The params.
    Definition: gemm_epilogue.h:57
    -
    cutlass::gemm::GemmEpilogue
    Definition: gemm_epilogue.h:53
    +
    cutlass::gemm::GemmEpilogue::Params
    Traits::Params Params
    The params.
    Definition: gemm_epilogue.h:46
    +
    cutlass::gemm::GemmEpilogue
    Definition: gemm_epilogue.h:42
    cutlass::gemm::IgemmGlobalLoadTransformer
    Definition: igemm_epilogue.h:167
    -
    cutlass::platform::is_same
    std::is_same (false specialization)
    Definition: platform.h:412
    +
    cutlass::platform::is_same
    std::is_same (false specialization)
    Definition: platform.h:420
    tile_iterator.h
    Defines the Tile Traits concept and iterators for loading and storing to tiles efficiently.
    cutlass::gemm::IgemmInt8ToFloatConverter::IgemmInt8ToFloatConverter
    CUTLASS_DEVICE IgemmInt8ToFloatConverter()
    Ctor.
    Definition: igemm_epilogue.h:123
    cutlass::gemm::IgemmEpilogueTraitsHelper::SharedStoreFragmentD
    SharedStoreIteratorD::Fragment SharedStoreFragmentD
    The fragment that needs to be passed to that store iterator.
    Definition: igemm_epilogue.h:229
    -
    cutlass::gemm::GemmEpilogueTraitsHelper< IgemmConfig_, EpilogueFunctor_, Index_ >::Scalar
    EpilogueFunctor_::Scalar Scalar
    The scalar.
    Definition: gemm_epilogue_traits.h:173
    +
    cutlass::gemm::GemmEpilogueTraitsHelper< IgemmConfig_, EpilogueFunctor_, Index_ >::Scalar
    EpilogueFunctor_::Scalar Scalar
    The scalar.
    Definition: gemm_epilogue_traits.h:188
    cutlass::gemm::IgemmEpilogueTraitsHelper
    Definition: igemm_epilogue.h:186
    -
    cutlass::MemorySpace::kGlobal
    Definition: load_store.h:43
    +
    cutlass::MemorySpace::kGlobal
    Definition: load_store.h:42
    cutlass::gemm::IgemmInt8ToFloatConverter::InputFragment
    Fragment< int8_t, kElements_ > InputFragment
    The input fragment.
    Definition: igemm_epilogue.h:115
    +
    cutlass::TileIteratorBase< TileTraits_, TileTraits_::Scalar, IteratorAdvance::kH, MemorySpace::kGlobal, Index_ >::Fragment
    Fragment< FragmentElement, ShapeCount< Iterations >::kCount *kAccessSize > Fragment
    The fragment.
    Definition: tile_iterator.h:196
    cutlass::gemm::IgemmEpilogue
    Definition: igemm_epilogue.h:290
    cutlass::gemm::IgemmFloatToInt8Converter
    Definition: igemm_epilogue.h:45
    cutlass::gemm::IgemmFloatToInt8Converter::transform
    CUTLASS_DEVICE void transform(Fragment_ const &src, int offset, OutputFragment &dst)
    Transform a fragment.
    Definition: igemm_epilogue.h:64
    -
    cutlass::gemm::GemmEpilogue::SharedStorage
    Traits::SharedStorage SharedStorage
    The shared storage.
    Definition: gemm_epilogue.h:59
    +
    cutlass::gemm::GemmEpilogue::SharedStorage
    Traits::SharedStorage SharedStorage
    The shared storage.
    Definition: gemm_epilogue.h:48
    cutlass::Fragment
    A template defining Fragment Concept.
    Definition: fragment.h:99
    -
    cutlass::IteratorAdvance::kH
    Definition: tile_iterator.h:62
    +
    cutlass::IteratorAdvance::kH
    Definition: tile_iterator.h:65
    cutlass::gemm::IgemmInt8ToFloatConverter::transform
    CUTLASS_DEVICE void transform(InputFragment const &src, OutputFragment &dst)
    Transform a fragment.
    Definition: igemm_epilogue.h:126
    cutlass::gemm::IgemmEpilogueTraitsHelper::Scalar
    Base::Scalar Scalar
    The scalar type of the epilogue.
    Definition: igemm_epilogue.h:194
    +
    cutlass::gemm::IgemmEpilogue::IgemmEpilogue
    CUTLASS_DEVICE IgemmEpilogue(typename Base::Params const &params_, typename Base::SharedStorage &shared_storage_, Coord< 3 > const &_problem_size)
    Ctor.
    Definition: igemm_epilogue.h:295
    cutlass::gemm::IgemmEpilogueTraitsHelper::GlobalFragmentC
    GlobalLoadIteratorC::Fragment GlobalFragmentC
    The fragment that needs to be produced by the load iterator.
    Definition: igemm_epilogue.h:205
    +
    cutlass::TileStoreIterator::Fragment
    Base::Fragment Fragment
    Fragment definition.
    Definition: tile_iterator.h:901
    cutlass::gemm::IgemmFloatToInt8Converter::transform
    CUTLASS_DEVICE void transform(InputFragment const &src, OutputFragment &dst)
    Transform a fragment.
    Definition: igemm_epilogue.h:58
    cutlass::gemm::IgemmFloatToInt8Converter::OutputFragment
    Fragment< int8_t, kElements_ > OutputFragment
    The output fragment.
    Definition: igemm_epilogue.h:49
    cutlass::gemm::IgemmEpilogueTraitsHelper::GlobalStoreIteratorD
    GemmGlobalIteratorCd< GlobalStoreTileTraits > GlobalStoreIteratorD
    The iterator to store to shared memory.
    Definition: igemm_epilogue.h:213
    cutlass::gemm::IgemmEpilogueTraitsHelper::SharedStoreTransformerD
    IgemmSharedStoreTransformer< typename IgemmConfig::Accumulators::Element, SharedStoreFragmentD >::Transformer SharedStoreTransformerD
    The transformer from accumulators to shared memory fragments.
    Definition: igemm_epilogue.h:233
    cutlass::gemm::IgemmEpilogueTraits::kInt8Output
    static bool const kInt8Output
    Do we output in int8?
    Definition: igemm_epilogue.h:283
    -
    cutlass::TileLoadIterator
    An iterator implementing Tile Load Iterator Concept for loading a tile from memory.
    Definition: tile_iterator.h:302
    +
    cutlass::TileLoadIterator
    An iterator implementing Tile Load Iterator Concept for loading a tile from memory.
    Definition: tile_iterator.h:399
    cutlass::gemm::IgemmSharedStoreTransformer::Transformer
    Convert< Fragment< InputScalar_, OutputFragment_::kElements >, OutputFragment_ > Transformer
    Definition: igemm_epilogue.h:180
    -
    cutlass::gemm::IgemmEpilogue< GemmEpilogueTraits_, true >::Base
    GemmEpilogue< GemmEpilogueTraits_ > Base
    The base class.
    Definition: igemm_epilogue.h:307
    +
    cutlass::gemm::IgemmEpilogue< GemmEpilogueTraits_, true >::Base
    GemmEpilogue< GemmEpilogueTraits_ > Base
    The base class.
    Definition: igemm_epilogue.h:306
    reshape_tile.h
    Defines a type for restructuring a tile.
    cutlass::gemm::IgemmEpilogueTraitsHelper::GlobalLoadTileTraits
    Base::GlobalLoadTileTraits GlobalLoadTileTraits
    The traits class for the iterator.
    Definition: igemm_epilogue.h:201
    cutlass::gemm::IgemmInt8ToFloatConverter::OutputFragment
    Fragment< float, kElements_ > OutputFragment
    The output fragment.
    Definition: igemm_epilogue.h:117
    cutlass::gemm::IgemmEpilogueTraitsHelper::Base
    GemmEpilogueTraitsHelper< IgemmConfig_, EpilogueFunctor_, Index_ > Base
    The base class.
    Definition: igemm_epilogue.h:189
    -
    cutlass::gemm::IgemmEpilogue::IgemmEpilogue
    CUTLASS_DEVICE IgemmEpilogue(typename Base::Params const &params_, typename Base::SharedStorage &shared_storage_, typename Base::Index m_, typename Base::Index n_)
    Ctor.
    Definition: igemm_epilogue.h:295
    -
    cutlass::gemm::GemmSharedLoadTileDTraits
    Definition: gemm_shared_tile.h:335
    -
    cutlass::gemm::GemmEpilogue::Index
    Traits::Index Index
    The index.
    Definition: gemm_epilogue.h:93
    +
    cutlass::gemm::GemmSharedLoadTileDTraits
    Definition: gemm_shared_tile.h:339
    cutlass::gemm::IgemmEpilogueTraitsHelper::GlobalFragmentD
    GlobalStoreIteratorD::Fragment GlobalFragmentD
    The fragment that needs to be passed to that store iterator.
    Definition: igemm_epilogue.h:215
    cutlass::gemm::IgemmEpilogueTraitsHelper::GlobalLoadIteratorC
    GemmGlobalIteratorCd< GlobalLoadTileTraits > GlobalLoadIteratorC
    The iterator to store to shared memory.
    Definition: igemm_epilogue.h:203
    -
    static_assert
    #define static_assert(__e, __m)
    Definition: platform.h:145
    +
    static_assert
    #define static_assert(__e, __m)
    Definition: platform.h:153
    cutlass::gemm::IgemmEpilogueTraitsHelper::IgemmConfig
    IgemmConfig_ IgemmConfig
    The config.
    Definition: igemm_epilogue.h:191
    -
    cutlass::gemm::IgemmEpilogue< GemmEpilogueTraits_, true >::IgemmEpilogue
    CUTLASS_DEVICE IgemmEpilogue(typename Base::Params const &params_, typename Base::SharedStorage &shared_storage_, typename Base::Index m_, typename Base::Index n_)
    Ctor.
    Definition: igemm_epilogue.h:310
    cutlass::Shape
    A Shape implementing Layout Concept describing the dimensions of a cube.
    Definition: shape.h:64
    cutlass::gemm::IgemmFloatToInt8Converter::IgemmFloatToInt8Converter
    CUTLASS_DEVICE IgemmFloatToInt8Converter()
    Ctor.
    Definition: igemm_epilogue.h:55
    cutlass::Fragment::Element
    Element_ Element
    The element.
    Definition: fragment.h:108
    cutlass::gemm::IgemmFloatToInt8Converter::InputFragment
    Fragment< float, kElements_ > InputFragment
    The input fragment.
    Definition: igemm_epilogue.h:47
    +
    cutlass::gemm::GemmEpilogueTraits
    Definition: gemm_epilogue_traits.h:70
    -
    cutlass::gemm::GemmGlobalIteratorCd
    Definition: gemm_global_tile.h:348
    +
    cutlass::gemm::GemmGlobalIteratorCd
    Definition: gemm_global_tile.h:396
    cutlass::gemm::IgemmSharedStoreTransformer
    Definition: igemm_epilogue.h:179
    gemm_global_stream.h
    Implements efficient loading of the thread block-level tile from global memory and storing to shared ...
    -
    cutlass::TileIteratorBase< TileTraits_, TileTraits_::Scalar, IteratorAdvance::kH, MemorySpace::kGlobal, Index_ >::Fragment
    Fragment< FragmentElement, ShapeCount< Iterations >::kCount *kAccessSize > Fragment
    The fragment.
    Definition: tile_iterator.h:154
    cutlass::Convert
    Definition: convert.h:38
    cutlass::gemm::IgemmGlobalStoreTransformer< float, Fragment< int8_t, kElements_ > >::Transformer
    IgemmFloatToInt8Converter< kElements_ > Transformer
    Definition: igemm_epilogue.h:107
    cutlass::gemm::IgemmEpilogueTraitsHelper::Iterations
    Base::Iterations Iterations
    The iterations.
    Definition: igemm_epilogue.h:196
    @@ -144,7 +143,7 @@ $(function() {
    cutlass::gemm::IgemmEpilogueTraitsHelper::GlobalStoreTileTraits
    Base::GlobalStoreTileTraits GlobalStoreTileTraits
    The traits class for the iterator.
    Definition: igemm_epilogue.h:211
    cutlass::gemm::IgemmGlobalLoadTransformer::Transformer
    Convert< InputFragment_, Fragment< OutputScalar_, InputFragment_::kElements > > Transformer
    Definition: igemm_epilogue.h:168
    fragment.h
    Defines Fragment, a statically-sized array for storing parts of matrices within a thread&#39;s registers...
    -
    cutlass::gemm::GemmSharedStoreTileDTraits::Scalar
    platform::remove_const< Scalar_ >::type Scalar
    The scalar.
    Definition: gemm_shared_tile.h:266
    +
    cutlass::gemm::GemmSharedStoreTileDTraits::Scalar
    platform::remove_const< Scalar_ >::type Scalar
    The scalar.
    Definition: gemm_shared_tile.h:272
    cutlass::gemm::IgemmInt8ToFloatConverter::transform
    CUTLASS_DEVICE void transform(Fragment_ const &src, int offset, OutputFragment &dst)
    Transform a fragment.
    Definition: igemm_epilogue.h:132
    cutlass::gemm::IgemmGlobalStoreTransformer::Transformer
    Convert< Fragment< InputScalar_, OutputFragment_::kElements >, OutputFragment_ > Transformer
    Definition: igemm_epilogue.h:102
    gemm_shared_stream.h
    Defines abstractions for managing loading and storing fragments to shared memory in the efficient GEM...
    @@ -153,14 +152,15 @@ $(function() {
    cutlass::gemm::IgemmGlobalLoadTransformer< Fragment< int8_t, kElements_ >, float >::Transformer
    IgemmInt8ToFloatConverter< kElements_ > Transformer
    Definition: igemm_epilogue.h:173
    convert.h
    Defines conversion operations among Fragments of different base type.
    cutlass::gemm::IgemmInt8ToFloatConverter
    Definition: igemm_epilogue.h:113
    -
    cutlass::gemm::GemmSharedLoadTileDTraits::Scalar
    platform::remove_const< Scalar_ >::type Scalar
    The scalar.
    Definition: gemm_shared_tile.h:337
    +
    cutlass::gemm::GemmSharedLoadTileDTraits::Scalar
    platform::remove_const< Scalar_ >::type Scalar
    The scalar.
    Definition: gemm_shared_tile.h:341
    +
    cutlass::gemm::IgemmEpilogue< GemmEpilogueTraits_, true >::IgemmEpilogue
    CUTLASS_DEVICE IgemmEpilogue(typename Base::Params const &params_, typename Base::SharedStorage &shared_storage_, Coord< 3 > const &_problem_size)
    Ctor.
    Definition: igemm_epilogue.h:309
    igemm_global_tile.h
    Implements tile iterators to partition the thread block tile into 2D subtiles and efficiently load ea...
    -
    cutlass::gemm::GemmSharedStoreTileDTraits
    Definition: gemm_shared_tile.h:264
    -
    cutlass::TileStoreIterator
    An iterator implementing Tile Store Iterator Concept for storing a tile to memory.
    Definition: tile_iterator.h:620
    +
    cutlass::gemm::GemmSharedStoreTileDTraits
    Definition: gemm_shared_tile.h:270
    +
    cutlass::TileStoreIterator
    An iterator implementing Tile Store Iterator Concept for storing a tile to memory.
    Definition: tile_iterator.h:836
    -Generated on Thu May 3 2018 16:36:12 for Cutlass by   +Generated on Wed Sep 19 2018 13:58:50 for Cutlass by   doxygen 1.8.14
    diff --git a/docs/igemm__global__tile_8h.html b/docs/igemm__global__tile_8h.html index d6a68016..4b5ee6d7 100644 --- a/docs/igemm__global__tile_8h.html +++ b/docs/igemm__global__tile_8h.html @@ -82,18 +82,20 @@ $(function() {

    Implements tile iterators to partition the thread block tile into 2D subtiles and efficiently load each. Applies permute transformation to construct 'interleaved K-strided' data layout in which 4-element dot products from the same K index are arranged in consecutive locations within shared memory. More...

    -
    #include <cutlass/coord.h>
    -#include <cutlass/gemm/gemm_global_tile.h>
    -#include <cutlass/matrix_traits.h>
    +
    #include "cutlass/coord.h"
    +#include "cutlass/gemm/gemm_global_tile.h"
    +#include "cutlass/matrix_traits.h"

    Go to the source code of this file.

    - + - - + + + +

    Classes

    struct  cutlass::gemm::IgemmContiguousGlobalTileTraits< kOperand_, kLayout_, Scalar_, Tile_, Threads_, kAccessSize_ >
    struct  cutlass::gemm::IgemmGlobalTileTraits< kOperand_, kLayout_, Scalar_, Tile_, Threads_, kAccessSize_ >
     
    struct  cutlass::gemm::IgemmContiguousGlobalTileTraits< kOperand_, kLayout_, Scalar_, Tile_, Threads_, kAccessSize_ >::ThreadOffset
     Computes the thread offset in (H, W) based on thread ID. More...
    struct  cutlass::gemm::IgemmGlobalTileTraits< kOperand_, kLayout_, Scalar_, Tile_, Threads_, kAccessSize_ >::ThreadOffset
     Computes the thread offset in (H, W) based on thread ID. More...
     
    struct  cutlass::gemm::IgemmGlobalIteratorAb< TileTraits_, Index_ >
     
    -Generated on Thu May 3 2018 16:36:12 for Cutlass by   +Generated on Wed Sep 19 2018 13:58:50 for Cutlass by   doxygen 1.8.14
    diff --git a/docs/igemm__global__tile_8h_source.html b/docs/igemm__global__tile_8h_source.html index df086169..04428a68 100644 --- a/docs/igemm__global__tile_8h_source.html +++ b/docs/igemm__global__tile_8h_source.html @@ -76,33 +76,46 @@ $(function() {
    igemm_global_tile.h
    -Go to the documentation of this file.
    1 /***************************************************************************************************
    2  * Copyright (c) 2017-2018, NVIDIA CORPORATION. All rights reserved.
    3  *
    4  * Redistribution and use in source and binary forms, with or without modification, are permitted
    5  * provided that the following conditions are met:
    6  * * Redistributions of source code must retain the above copyright notice, this list of
    7  * conditions and the following disclaimer.
    8  * * Redistributions in binary form must reproduce the above copyright notice, this list of
    9  * conditions and the following disclaimer in the documentation and/or other materials
    10  * provided with the distribution.
    11  * * Neither the name of the NVIDIA CORPORATION nor the names of its contributors may be used
    12  * to endorse or promote products derived from this software without specific prior written
    13  * permission.
    14  *
    15  * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND ANY EXPRESS OR
    16  * IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND
    17  * FITNESS FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE
    18  * FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING,
    19  * BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS;
    20  * OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT,
    21  * STRICT LIABILITY, OR TOR (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
    22  * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
    23  *
    24  **************************************************************************************************/
    33 #pragma once
    34 
    35 #include <cutlass/coord.h>
    36 #include <cutlass/gemm/gemm_global_tile.h>
    37 #include <cutlass/matrix_traits.h>
    38 
    39 namespace cutlass {
    40 namespace gemm {
    41 
    43 
    44 template <GemmOperand::Kind kOperand_,
    45  MatrixLayout::Kind kLayout_,
    46  typename Scalar_,
    47  typename Tile_,
    48  typename Threads_,
    49  int kAccessSize_>
    50 struct IgemmContiguousGlobalTileTraits : public GemmGlobalTileTraits<
    51  // Which GEMM operand?
    52  kOperand_,
    53  // The layout.
    54  kLayout_,
    55  // The scalar.
    56  Scalar_,
    57  // The tile.
    58  Tile_,
    59  // The threads.
    60  Threads_,
    61  // The number of scalars per LDG/STG.
    62  kAccessSize_> {
    64  typedef GemmGlobalTileTraits<kOperand_, kLayout_, Scalar_, Tile_, Threads_, kAccessSize_> Base;
    66  typedef typename Base::Threads Threads;
    68  typedef Shape<Base::Threads::kH * 4, 1, Base::Threads::kW, Base::kAccessSize> Delta;
    70  typedef Shape<Base::Tile::kH / Base::Threads::kH / 4,
    71  4,
    72  Base::Tile::kW / Base::Threads::kW,
    73  Base::Tile::kC / Base::kAccessSize>
    74  Iterations;
    75 
    77  struct ThreadOffset {
    78  CUTLASS_HOST_DEVICE
    79  Coord<4> operator()() const {
    80  int thread_offset_h = threadIdx.x / Threads::kW * ThreadsDelta::kH;
    81  int thread_offset_w = threadIdx.x % Threads::kW * ThreadsDelta::kW;
    82 
    83  return make_Coord(0, thread_offset_h, thread_offset_w, 0);
    84  }
    85  };
    86 
    87  public:
    89  typedef Shape<1, 4, Base::Tile::kC> ThreadsDelta;
    90 };
    91 
    93 
    94 } // namespace gemm
    95 } // namespace cutlass
    cutlass::gemm::IgemmContiguousGlobalTileTraits::ThreadOffset
    Computes the thread offset in (H, W) based on thread ID.
    Definition: igemm_global_tile.h:77
    -
    cutlass
    Definition: convert.h:33
    +Go to the documentation of this file.
    1 /***************************************************************************************************
    2  * Copyright (c) 2017-2018, NVIDIA CORPORATION. All rights reserved.
    3  *
    4  * Redistribution and use in source and binary forms, with or without modification, are permitted
    5  * provided that the following conditions are met:
    6  * * Redistributions of source code must retain the above copyright notice, this list of
    7  * conditions and the following disclaimer.
    8  * * Redistributions in binary form must reproduce the above copyright notice, this list of
    9  * conditions and the following disclaimer in the documentation and/or other materials
    10  * provided with the distribution.
    11  * * Neither the name of the NVIDIA CORPORATION nor the names of its contributors may be used
    12  * to endorse or promote products derived from this software without specific prior written
    13  * permission.
    14  *
    15  * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND ANY EXPRESS OR
    16  * IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND
    17  * FITNESS FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE
    18  * FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING,
    19  * BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS;
    20  * OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT,
    21  * STRICT LIABILITY, OR TOR (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
    22  * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
    23  *
    24  **************************************************************************************************/
    33 #pragma once
    34 
    35 #include "cutlass/coord.h"
    36 #include "cutlass/gemm/gemm_global_tile.h"
    37 #include "cutlass/matrix_traits.h"
    38 
    39 namespace cutlass {
    40 namespace gemm {
    41 
    43 
    44 template <GemmOperand::Kind kOperand_,
    45  MatrixLayout::Kind kLayout_,
    46  typename Scalar_,
    47  typename Tile_,
    48  typename Threads_,
    49  int kAccessSize_>
    50 struct IgemmGlobalTileTraits : public GemmGlobalTileTraits<
    51  // Which GEMM operand?
    52  kOperand_,
    53  // The layout.
    54  kLayout_,
    55  // The scalar.
    56  Scalar_,
    57  // The tile.
    58  Tile_,
    59  // The threads.
    60  Threads_,
    61  // The number of scalars per LDG/STG.
    62  kAccessSize_> {
    64  typedef GemmGlobalTileTraits<kOperand_, kLayout_, Scalar_, Tile_, Threads_, kAccessSize_> Base;
    66  typedef typename Base::Threads Threads;
    68  typedef Shape<Base::Threads::kH * 4, 1, Base::Threads::kW, Base::kAccessSize> Delta;
    70  typedef Shape<Base::VectorizedTile::kH / Base::Threads::kH / 4,
    71  4,
    72  Base::VectorizedTile::kW / Base::Threads::kW,
    73  Base::VectorizedTile::kC / Base::kAccessSize>
    74  Iterations;
    75 
    77  struct ThreadOffset {
    78  CUTLASS_HOST_DEVICE
    79  Coord<4> operator()() const {
    80  int thread_offset_h = threadIdx.x / Threads::kW * ThreadsDelta::kH;
    81  int thread_offset_w = threadIdx.x % Threads::kW * ThreadsDelta::kW;
    82 
    83  return make_Coord(0, thread_offset_h, thread_offset_w, 0);
    84  }
    85  };
    86 
    87  public:
    89  typedef Shape<1, 4, Base::VectorizedTile::kC> ThreadsDelta;
    90 };
    91 
    93 
    94 template <typename TileTraits_, typename Index_ = int>
    95 struct IgemmGlobalIteratorAb : public GemmGlobalIteratorAb<TileTraits_, Index_> {
    97  typedef GemmGlobalIteratorAb<TileTraits_, Index_> Base;
    99  typedef typename TileTraits_::ThreadOffset ThreadOffset;
    100 
    102  CUTLASS_DEVICE IgemmGlobalIteratorAb(typename Base::Params const& _params,
    103  const Coord<3>& bounds,
    104  const Coord<3>& threadblock_offset,
    105  ThreadOffset thread_offset_func = ThreadOffset())
    106  : Base(_params, bounds, threadblock_offset, thread_offset_func), mask_(0xffffffff) {
    107  // The number of elements read in a single iteration.
    108  int const kBlock = TileTraits_::Tile::kW;
    109  // The residue.
    110  int const kResidue = (int)(bounds[1] % kBlock);
    111 
    112  // Compute the number of elements that are valid.
    113  int const left = kResidue - Base::thread_offset[2];
    114  if (left > 0 && left < 4) {
    115  mask_ = (1u << (8 * left)) - 1u;
    116  }
    117  }
    118 
    119  CUTLASS_DEVICE void load_element(
    120  typename Base::AccessType& value, int d, int h, int w, int c) const {
    121  Base::load_element(value, d, h, w, c);
    122  reinterpret_cast<uint32_t&>(value) &= mask_;
    123  }
    124 
    126  uint32_t mask_;
    127 };
    128 
    130 
    131 } // namespace gemm
    132 } // namespace cutlass
    cutlass
    Definition: convert.h:33
    +
    cutlass::gemm::IgemmGlobalTileTraits::Threads
    Base::Threads Threads
    The threads.
    Definition: igemm_global_tile.h:66
    +
    cutlass::gemm::IgemmGlobalTileTraits::ThreadOffset
    Computes the thread offset in (H, W) based on thread ID.
    Definition: igemm_global_tile.h:77
    gemm_global_tile.h
    Defines iterators for efficiently loading and storing to global memory.
    cutlass::gemm::GemmGlobalTileTraits
    Definition: gemm_global_tile.h:70
    coord.h
    A Coord is a coordinate of arbitrary rank into a tensor or matrix.
    -
    cutlass::make_Coord
    CUTLASS_HOST_DEVICE Coord< 1 > make_Coord(int _0)
    Helper to make a 2-element coordinate.
    Definition: coord.h:241
    -
    cutlass::gemm::IgemmContiguousGlobalTileTraits::Delta
    Shape< Base::Threads::kH *4, 1, Base::Threads::kW, Base::kAccessSize > Delta
    The strides in each dimension between different loads/stores.
    Definition: igemm_global_tile.h:68
    +
    cutlass::gemm::IgemmGlobalTileTraits::Iterations
    Shape< Base::VectorizedTile::kH/Base::Threads::kH/4, 4, Base::VectorizedTile::kW/Base::Threads::kW, Base::VectorizedTile::kC/Base::kAccessSize > Iterations
    The number of iterations needed to load/store the tile.
    Definition: igemm_global_tile.h:74
    +
    cutlass::make_Coord
    CUTLASS_HOST_DEVICE Coord< 1 > make_Coord(int _0)
    Helper to make a 2-element coordinate.
    Definition: coord.h:318
    +
    cutlass::gemm::GemmGlobalIteratorAb::load_element
    CUTLASS_HOST_DEVICE void load_element(typename Base::AccessType &value, int d, int h, int w, int c) const
    Loads a single fragment element from memory.
    Definition: gemm_global_tile.h:292
    +
    cutlass::gemm::IgemmGlobalTileTraits::ThreadOffset::operator()
    CUTLASS_HOST_DEVICE Coord< 4 > operator()() const
    Definition: igemm_global_tile.h:79
    +
    cutlass::gemm::GemmGlobalIteratorAb
    Definition: gemm_global_tile.h:163
    cutlass::Shape::kH
    static int const kH
    The height of the cube.
    Definition: shape.h:68
    -
    cutlass::gemm::IgemmContiguousGlobalTileTraits::Base
    GemmGlobalTileTraits< kOperand_, kLayout_, Scalar_, Tile_, Threads_, kAccessSize_ > Base
    The base class.
    Definition: igemm_global_tile.h:64
    -
    cutlass::gemm::IgemmContiguousGlobalTileTraits::Iterations
    Shape< Base::Tile::kH/Base::Threads::kH/4, 4, Base::Tile::kW/Base::Threads::kW, Base::Tile::kC/Base::kAccessSize > Iterations
    The number of iterations needed to load/store the tile.
    Definition: igemm_global_tile.h:74
    +
    cutlass::TileLoadIterator
    An iterator implementing Tile Load Iterator Concept for loading a tile from memory.
    Definition: tile_iterator.h:399
    +
    cutlass::gemm::IgemmGlobalTileTraits
    Definition: igemm_global_tile.h:50
    +
    cutlass::gemm::IgemmGlobalIteratorAb::load_element
    CUTLASS_DEVICE void load_element(typename Base::AccessType &value, int d, int h, int w, int c) const
    Definition: igemm_global_tile.h:119
    +
    cutlass::gemm::IgemmGlobalIteratorAb::Base
    GemmGlobalIteratorAb< TileTraits_, Index_ > Base
    The base class.
    Definition: igemm_global_tile.h:97
    +
    cutlass::gemm::IgemmGlobalIteratorAb
    Definition: igemm_global_tile.h:95
    CUTLASS_HOST_DEVICE
    #define CUTLASS_HOST_DEVICE
    Definition: cutlass.h:46
    -
    cutlass::gemm::IgemmContiguousGlobalTileTraits
    Definition: igemm_global_tile.h:50
    +
    cutlass::Vector
    Definition: vector.h:62
    cutlass::Shape
    A Shape implementing Layout Concept describing the dimensions of a cube.
    Definition: shape.h:64
    +
    cutlass::gemm::IgemmGlobalIteratorAb::ThreadOffset
    TileTraits_::ThreadOffset ThreadOffset
    The functor to compute the thread offset.
    Definition: igemm_global_tile.h:99
    +
    cutlass::gemm::IgemmGlobalIteratorAb::mask_
    uint32_t mask_
    The mask to clean up the values.
    Definition: igemm_global_tile.h:126
    +
    cutlass::gemm::GemmGlobalTileTraits::Threads
    ReshapeThreads< VectorizedTile, Threads_ >::Threads Threads
    The threads shape.
    Definition: gemm_global_tile.h:88
    +
    cutlass::gemm::IgemmGlobalIteratorAb::IgemmGlobalIteratorAb
    CUTLASS_DEVICE IgemmGlobalIteratorAb(typename Base::Params const &_params, const Coord< 3 > &bounds, const Coord< 3 > &threadblock_offset, ThreadOffset thread_offset_func=ThreadOffset())
    Constructor.
    Definition: igemm_global_tile.h:102
    +
    cutlass::gemm::IgemmGlobalTileTraits::ThreadsDelta
    Shape< 1, 4, Base::VectorizedTile::kC > ThreadsDelta
    The threads strides.
    Definition: igemm_global_tile.h:89
    +
    cutlass::gemm::GemmGlobalIteratorAb::ThreadOffset
    TileTraits_::ThreadOffset ThreadOffset
    The thread offset.
    Definition: gemm_global_tile.h:192
    cutlass::Shape::kW
    static int const kW
    The width of the cube.
    Definition: shape.h:70
    -
    cutlass::MatrixLayout::Kind
    Kind
    Definition: matrix_traits.h:36
    +
    cutlass::TileLoadIterator::Params
    Parameters.
    Definition: tile_iterator.h:491
    +
    cutlass::MatrixLayout::Kind
    Kind
    Enumeration defining fundamental contiguous layouts.
    Definition: matrix_traits.h:159
    cutlass::gemm::GemmGlobalTileTraits::kAccessSize
    static int const kAccessSize
    The number of scalars per LDG/STG.
    Definition: gemm_global_tile.h:80
    -
    cutlass::GemmOperand::Kind
    Kind
    Definition: matrix_traits.h:43
    -
    cutlass::gemm::GemmGlobalTileTraits::Threads
    ReshapeThreads< Tile, Threads_ >::Threads Threads
    The threads shape.
    Definition: gemm_global_tile.h:87
    +
    cutlass::GemmOperand::Kind
    Kind
    Definition: matrix_traits.h:357
    +
    cutlass::gemm::IgemmGlobalTileTraits::Delta
    Shape< Base::Threads::kH *4, 1, Base::Threads::kW, Base::kAccessSize > Delta
    The strides in each dimension between different loads/stores.
    Definition: igemm_global_tile.h:68
    matrix_traits.h
    Defines properties of matrices used to denote layout and operands to GEMM kernels.
    -
    cutlass::gemm::IgemmContiguousGlobalTileTraits::ThreadsDelta
    Shape< 1, 4, Base::Tile::kC > ThreadsDelta
    The threads strides.
    Definition: igemm_global_tile.h:89
    -
    cutlass::gemm::IgemmContiguousGlobalTileTraits::ThreadOffset::operator()
    CUTLASS_HOST_DEVICE Coord< 4 > operator()() const
    Definition: igemm_global_tile.h:79
    -
    cutlass::gemm::IgemmContiguousGlobalTileTraits::Threads
    Base::Threads Threads
    The threads.
    Definition: igemm_global_tile.h:66
    +
    cutlass::gemm::GemmGlobalIteratorAb::thread_offset
    Coord< 4 > thread_offset
    Offset of an individual lane from the start of the tile.
    Definition: gemm_global_tile.h:237
    +
    cutlass::gemm::IgemmGlobalTileTraits::Base
    GemmGlobalTileTraits< kOperand_, kLayout_, Scalar_, Tile_, Threads_, kAccessSize_ > Base
    The base class.
    Definition: igemm_global_tile.h:64
    -Generated on Thu May 3 2018 16:36:12 for Cutlass by   +Generated on Wed Sep 19 2018 13:58:50 for Cutlass by   doxygen 1.8.14
    diff --git a/docs/igemm__multiply__add_8h.html b/docs/igemm__multiply__add_8h.html index 266cb5f1..d67e57b8 100644 --- a/docs/igemm__multiply__add_8h.html +++ b/docs/igemm__multiply__add_8h.html @@ -82,15 +82,15 @@ $(function() {

    Implements matrix multiply accumulate operation of 8-bit integer data using DP4A instruction. More...

    -
    #include <cutlass/fragment.h>
    -#include <cutlass/gemm/thread_multiply_add.h>
    +
    #include "cutlass/fragment.h"
    +#include "cutlass/gemm/thread_multiply_add.h"

    Go to the source code of this file.

    @@ -108,7 +110,7 @@ Namespaces

    - - + +

    Classes

    struct  cutlass::gemm::ThreadMultiplyAdd< AccumulatorsPerThread_, ThreadsPerWarp_, int8_t, int8_t, int >
     Template performing matrix multiply-add operation within a thread. More...
    struct  cutlass::gemm::ThreadMultiplyAdd< ThreadGemmShape_, ThreadsPerWarp_, int8_t, int8_t, int >
     Template performing matrix multiply-add operation within a thread. More...
     
    -Generated on Thu May 3 2018 16:36:12 for Cutlass by   +Generated on Wed Sep 19 2018 13:58:50 for Cutlass by   doxygen 1.8.14
    diff --git a/docs/igemm__multiply__add_8h_source.html b/docs/igemm__multiply__add_8h_source.html index 414c2ce1..b67129ef 100644 --- a/docs/igemm__multiply__add_8h_source.html +++ b/docs/igemm__multiply__add_8h_source.html @@ -76,29 +76,30 @@ $(function() {
    igemm_multiply_add.h
    -Go to the documentation of this file.
    1 /***************************************************************************************************
    2  * Copyright (c) 2017-2018, NVIDIA CORPORATION. All rights reserved.
    3  *
    4  * Redistribution and use in source and binary forms, with or without modification, are permitted
    5  * provided that the following conditions are met:
    6  * * Redistributions of source code must retain the above copyright notice, this list of
    7  * conditions and the following disclaimer.
    8  * * Redistributions in binary form must reproduce the above copyright notice, this list of
    9  * conditions and the following disclaimer in the documentation and/or other materials
    10  * provided with the distribution.
    11  * * Neither the name of the NVIDIA CORPORATION nor the names of its contributors may be used
    12  * to endorse or promote products derived from this software without specific prior written
    13  * permission.
    14  *
    15  * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND ANY EXPRESS OR
    16  * IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND
    17  * FITNESS FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE
    18  * FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING,
    19  * BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS;
    20  * OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT,
    21  * STRICT LIABILITY, OR TOR (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
    22  * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
    23  *
    24  **************************************************************************************************/
    29 #pragma once
    30 
    31 #include <cutlass/fragment.h>
    32 
    33 #include <cutlass/gemm/thread_multiply_add.h>
    34 
    35 namespace cutlass {
    36 namespace gemm {
    37 
    39 
    41 template <typename AccumulatorsPerThread_, typename ThreadsPerWarp_>
    42 struct ThreadMultiplyAdd<AccumulatorsPerThread_, ThreadsPerWarp_, int8_t, int8_t, int> {
    44  typedef Shape<4, 1, 1> InstructionShape;
    46  typedef AccumulatorsPerThread_ AccumulatorsPerThread;
    48  typedef ThreadsPerWarp_ ThreadsPerWarp;
    50  typedef typename ShapeMul<AccumulatorsPerThread, ThreadsPerWarp>::Shape AccumulatorsPerWarp;
    52  typedef int8_t ScalarA;
    54  typedef Fragment<ScalarA, AccumulatorsPerThread::kW * 4> FragmentA;
    56  typedef int8_t ScalarB;
    58  typedef Fragment<ScalarB, AccumulatorsPerThread::kH * 4> FragmentB;
    60  typedef int ScalarC;
    62  typedef Fragment<ScalarC, AccumulatorsPerThread::kH * AccumulatorsPerThread::kW> Accumulators;
    63 
    65  CUTLASS_DEVICE ThreadMultiplyAdd() {}
    66 
    68  CUTLASS_DEVICE void multiply_add(FragmentA const& a,
    69  FragmentB const& b,
    70  Accumulators const& c,
    71  Accumulators& d) {
    72  // The inputs.
    73  int const* a_int = reinterpret_cast<int const*>(&a[0]);
    74  int const* b_int = reinterpret_cast<int const*>(&b[0]);
    75 
    76  for (int j = 0; j < AccumulatorsPerThread::kH; ++j) {
    77  for (int i = 0; i < AccumulatorsPerThread::kW; ++i) {
    78  asm volatile("dp4a.s32.s32 %0, %1, %2, %3;"
    79  : "=r"(d[j * AccumulatorsPerThread::kW + i])
    80  : "r"(a_int[i]), "r"(b_int[j]), "r"(c[j * AccumulatorsPerThread::kW + i]));
    81  }
    82  }
    83  }
    84 };
    85 
    87 
    88 } // namespace gemm
    89 } // namespace cutlass
    cutlass::gemm::ThreadMultiplyAdd< AccumulatorsPerThread_, ThreadsPerWarp_, int8_t, int8_t, int >::ScalarA
    int8_t ScalarA
    The type for A.
    Definition: igemm_multiply_add.h:52
    -
    cutlass
    Definition: convert.h:33
    +Go to the documentation of this file.
    1 /***************************************************************************************************
    2  * Copyright (c) 2017-2018, NVIDIA CORPORATION. All rights reserved.
    3  *
    4  * Redistribution and use in source and binary forms, with or without modification, are permitted
    5  * provided that the following conditions are met:
    6  * * Redistributions of source code must retain the above copyright notice, this list of
    7  * conditions and the following disclaimer.
    8  * * Redistributions in binary form must reproduce the above copyright notice, this list of
    9  * conditions and the following disclaimer in the documentation and/or other materials
    10  * provided with the distribution.
    11  * * Neither the name of the NVIDIA CORPORATION nor the names of its contributors may be used
    12  * to endorse or promote products derived from this software without specific prior written
    13  * permission.
    14  *
    15  * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND ANY EXPRESS OR
    16  * IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND
    17  * FITNESS FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE
    18  * FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING,
    19  * BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS;
    20  * OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT,
    21  * STRICT LIABILITY, OR TOR (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
    22  * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
    23  *
    24  **************************************************************************************************/
    29 #pragma once
    30 
    31 #include "cutlass/fragment.h"
    32 
    33 #include "cutlass/gemm/thread_multiply_add.h"
    34 
    35 namespace cutlass {
    36 namespace gemm {
    37 
    39 
    41 template <typename ThreadGemmShape_, typename ThreadsPerWarp_>
    42 struct ThreadMultiplyAdd<ThreadGemmShape_, ThreadsPerWarp_, int8_t, int8_t, int> {
    44  typedef Shape<4, 1, 1> InstructionShape;
    46  typedef ThreadGemmShape_ ThreadGemmShape;
    48  typedef ThreadGemmShape AccumulatorsPerThread;
    50  typedef ThreadsPerWarp_ ThreadsPerWarp;
    52  typedef typename ShapeMul<ThreadGemmShape, ThreadsPerWarp>::Shape AccumulatorsPerWarp;
    54  typedef int8_t ScalarA;
    56  typedef Fragment<ScalarA, AccumulatorsPerThread::kW * 4> FragmentA;
    58  typedef int8_t ScalarB;
    60  typedef Fragment<ScalarB, AccumulatorsPerThread::kH * 4> FragmentB;
    62  typedef int ScalarC;
    64  typedef Fragment<ScalarC, AccumulatorsPerThread::kH * AccumulatorsPerThread::kW> Accumulators;
    65 
    67  CUTLASS_DEVICE ThreadMultiplyAdd() {}
    68 
    70  CUTLASS_DEVICE void multiply_add(FragmentA const& a,
    71  FragmentB const& b,
    72  Accumulators const& c,
    73  Accumulators& d) {
    74  // The inputs.
    75  int const* a_int = reinterpret_cast<int const*>(&a[0]);
    76  int const* b_int = reinterpret_cast<int const*>(&b[0]);
    77 
    78  for (int j = 0; j < AccumulatorsPerThread::kH; ++j) {
    79  for (int i = 0; i < AccumulatorsPerThread::kW; ++i) {
    80  asm volatile("dp4a.s32.s32 %0, %1, %2, %3;"
    81  : "=r"(d[j * AccumulatorsPerThread::kW + i])
    82  : "r"(a_int[i]), "r"(b_int[j]), "r"(c[j * AccumulatorsPerThread::kW + i]));
    83  }
    84  }
    85  }
    86 };
    87 
    89 
    90 } // namespace gemm
    91 } // namespace cutlass
    cutlass
    Definition: convert.h:33
    +
    cutlass::gemm::ThreadMultiplyAdd< ThreadGemmShape_, ThreadsPerWarp_, int8_t, int8_t, int >::FragmentA
    Fragment< ScalarA, AccumulatorsPerThread::kW *4 > FragmentA
    The fragment for A.
    Definition: igemm_multiply_add.h:56
    cutlass::ShapeMul::Shape
    Shape< A_::kD *B_::kD, A_::kH *B_::kH, A_::kW *B_::kW, A_::kC *B_::kC > Shape
    Definition: shape.h:119
    cutlass::Fragment
    A template defining Fragment Concept.
    Definition: fragment.h:99
    thread_multiply_add.h
    Template implementing matrix multiply-add operations on fragments.
    -
    cutlass::gemm::ThreadMultiplyAdd< AccumulatorsPerThread_, ThreadsPerWarp_, int8_t, int8_t, int >::Accumulators
    Fragment< ScalarC, AccumulatorsPerThread::kH *AccumulatorsPerThread::kW > Accumulators
    The accumulators.
    Definition: igemm_multiply_add.h:62
    -
    cutlass::gemm::ThreadMultiplyAdd< AccumulatorsPerThread_, ThreadsPerWarp_, int8_t, int8_t, int >::AccumulatorsPerWarp
    ShapeMul< AccumulatorsPerThread, ThreadsPerWarp >::Shape AccumulatorsPerWarp
    The number of accumulators per warp.
    Definition: igemm_multiply_add.h:50
    -
    cutlass::gemm::ThreadMultiplyAdd< AccumulatorsPerThread_, ThreadsPerWarp_, int8_t, int8_t, int >::FragmentB
    Fragment< ScalarB, AccumulatorsPerThread::kH *4 > FragmentB
    The fragment for B.
    Definition: igemm_multiply_add.h:58
    -
    cutlass::gemm::ThreadMultiplyAdd< AccumulatorsPerThread_, ThreadsPerWarp_, int8_t, int8_t, int >::ThreadMultiplyAdd
    CUTLASS_DEVICE ThreadMultiplyAdd()
    Ctor.
    Definition: igemm_multiply_add.h:65
    -
    cutlass::gemm::ThreadMultiplyAdd< AccumulatorsPerThread_, ThreadsPerWarp_, int8_t, int8_t, int >::InstructionShape
    Shape< 4, 1, 1 > InstructionShape
    The shape of the instruction.
    Definition: igemm_multiply_add.h:44
    -
    cutlass::gemm::ThreadMultiplyAdd< AccumulatorsPerThread_, ThreadsPerWarp_, int8_t, int8_t, int >::ThreadsPerWarp
    ThreadsPerWarp_ ThreadsPerWarp
    The number of threads per warp.
    Definition: igemm_multiply_add.h:48
    -
    cutlass::gemm::ThreadMultiplyAdd< AccumulatorsPerThread_, ThreadsPerWarp_, int8_t, int8_t, int >::AccumulatorsPerThread
    AccumulatorsPerThread_ AccumulatorsPerThread
    The number of accumulators per thread.
    Definition: igemm_multiply_add.h:46
    +
    cutlass::gemm::ThreadMultiplyAdd< ThreadGemmShape_, ThreadsPerWarp_, int8_t, int8_t, int >::ThreadMultiplyAdd
    CUTLASS_DEVICE ThreadMultiplyAdd()
    Ctor.
    Definition: igemm_multiply_add.h:67
    +
    cutlass::gemm::ThreadMultiplyAdd< ThreadGemmShape_, ThreadsPerWarp_, int8_t, int8_t, int >::multiply_add
    CUTLASS_DEVICE void multiply_add(FragmentA const &a, FragmentB const &b, Accumulators const &c, Accumulators &d)
    Multiply : d = a*b + c.
    Definition: igemm_multiply_add.h:70
    +
    cutlass::gemm::ThreadMultiplyAdd< ThreadGemmShape_, ThreadsPerWarp_, int8_t, int8_t, int >::ScalarC
    int ScalarC
    The type for C and D.
    Definition: igemm_multiply_add.h:62
    +
    cutlass::gemm::ThreadMultiplyAdd< ThreadGemmShape_, ThreadsPerWarp_, int8_t, int8_t, int >::InstructionShape
    Shape< 4, 1, 1 > InstructionShape
    The shape of the instruction.
    Definition: igemm_multiply_add.h:44
    +
    cutlass::gemm::ThreadMultiplyAdd< ThreadGemmShape_, ThreadsPerWarp_, int8_t, int8_t, int >::ThreadsPerWarp
    ThreadsPerWarp_ ThreadsPerWarp
    The number of threads per warp.
    Definition: igemm_multiply_add.h:50
    +
    cutlass::gemm::ThreadMultiplyAdd< ThreadGemmShape_, ThreadsPerWarp_, int8_t, int8_t, int >::ThreadGemmShape
    ThreadGemmShape_ ThreadGemmShape
    Shape of the thread-level GEMM (K-by-N-by-M)
    Definition: igemm_multiply_add.h:46
    +
    cutlass::gemm::ThreadMultiplyAdd< ThreadGemmShape_, ThreadsPerWarp_, int8_t, int8_t, int >::Accumulators
    Fragment< ScalarC, AccumulatorsPerThread::kH *AccumulatorsPerThread::kW > Accumulators
    The accumulators.
    Definition: igemm_multiply_add.h:64
    cutlass::Shape
    A Shape implementing Layout Concept describing the dimensions of a cube.
    Definition: shape.h:64
    -
    cutlass::gemm::ThreadMultiplyAdd
    Template performing matrix multiply-add operation within a thread.
    Definition: thread_multiply_add.h:43
    -
    cutlass::gemm::ThreadMultiplyAdd< AccumulatorsPerThread_, ThreadsPerWarp_, int8_t, int8_t, int >::FragmentA
    Fragment< ScalarA, AccumulatorsPerThread::kW *4 > FragmentA
    The fragment for A.
    Definition: igemm_multiply_add.h:54
    -
    cutlass::gemm::ThreadMultiplyAdd< AccumulatorsPerThread_, ThreadsPerWarp_, int8_t, int8_t, int >::ScalarC
    int ScalarC
    The type for C and D.
    Definition: igemm_multiply_add.h:60
    -
    cutlass::gemm::ThreadMultiplyAdd< AccumulatorsPerThread_, ThreadsPerWarp_, int8_t, int8_t, int >::ScalarB
    int8_t ScalarB
    The type for B.
    Definition: igemm_multiply_add.h:56
    -
    cutlass::gemm::ThreadMultiplyAdd< AccumulatorsPerThread_, ThreadsPerWarp_, int8_t, int8_t, int >::multiply_add
    CUTLASS_DEVICE void multiply_add(FragmentA const &a, FragmentB const &b, Accumulators const &c, Accumulators &d)
    Multiply : d = a*b + c.
    Definition: igemm_multiply_add.h:68
    +
    cutlass::gemm::ThreadMultiplyAdd< ThreadGemmShape_, ThreadsPerWarp_, int8_t, int8_t, int >::AccumulatorsPerWarp
    ShapeMul< ThreadGemmShape, ThreadsPerWarp >::Shape AccumulatorsPerWarp
    The number of accumulators per warp.
    Definition: igemm_multiply_add.h:52
    +
    cutlass::gemm::ThreadMultiplyAdd
    Template performing matrix multiply-add operation within a thread.
    Definition: thread_multiply_add.h:44
    +
    cutlass::gemm::ThreadMultiplyAdd< ThreadGemmShape_, ThreadsPerWarp_, int8_t, int8_t, int >::AccumulatorsPerThread
    ThreadGemmShape AccumulatorsPerThread
    Aliased for compatibility. Will be removed in CUTLASS v2.0.
    Definition: igemm_multiply_add.h:48
    +
    cutlass::gemm::ThreadMultiplyAdd< ThreadGemmShape_, ThreadsPerWarp_, int8_t, int8_t, int >::FragmentB
    Fragment< ScalarB, AccumulatorsPerThread::kH *4 > FragmentB
    The fragment for B.
    Definition: igemm_multiply_add.h:60
    +
    cutlass::gemm::ThreadMultiplyAdd< ThreadGemmShape_, ThreadsPerWarp_, int8_t, int8_t, int >::ScalarA
    int8_t ScalarA
    The type for A.
    Definition: igemm_multiply_add.h:54
    fragment.h
    Defines Fragment, a statically-sized array for storing parts of matrices within a thread&#39;s registers...
    +
    cutlass::gemm::ThreadMultiplyAdd< ThreadGemmShape_, ThreadsPerWarp_, int8_t, int8_t, int >::ScalarB
    int8_t ScalarB
    The type for B.
    Definition: igemm_multiply_add.h:58
    -Generated on Thu May 3 2018 16:36:12 for Cutlass by   +Generated on Wed Sep 19 2018 13:58:50 for Cutlass by   doxygen 1.8.14
    diff --git a/docs/igemm__swizzle_8h.html b/docs/igemm__swizzle_8h.html index a631d215..c8785521 100644 --- a/docs/igemm__swizzle_8h.html +++ b/docs/igemm__swizzle_8h.html @@ -82,7 +82,7 @@ $(function() {

    Transposes a fragment of data containing packed 8-bit integer elements. More...

    -
    #include <cutlass/fragment.h>
    +
    #include "cutlass/fragment.h"

    Go to the source code of this file.

    @@ -103,7 +103,7 @@ Namespaces

    @@ -101,7 +101,7 @@ Namespaces
    -Generated on Thu May 3 2018 16:36:12 for Cutlass by   +Generated on Wed Sep 19 2018 13:58:50 for Cutlass by   doxygen 1.8.14
    diff --git a/docs/igemm__swizzle_8h_source.html b/docs/igemm__swizzle_8h_source.html index 93990830..015b5f9a 100644 --- a/docs/igemm__swizzle_8h_source.html +++ b/docs/igemm__swizzle_8h_source.html @@ -76,14 +76,14 @@ $(function() {
    igemm_swizzle.h
    -Go to the documentation of this file.
    1 /***************************************************************************************************
    2  * Copyright (c) 2017-2018, NVIDIA CORPORATION. All rights reserved.
    3  *
    4  * Redistribution and use in source and binary forms, with or without modification, are permitted
    5  * provided that the following conditions are met:
    6  * * Redistributions of source code must retain the above copyright notice, this list of
    7  * conditions and the following disclaimer.
    8  * * Redistributions in binary form must reproduce the above copyright notice, this list of
    9  * conditions and the following disclaimer in the documentation and/or other materials
    10  * provided with the distribution.
    11  * * Neither the name of the NVIDIA CORPORATION nor the names of its contributors may be used
    12  * to endorse or promote products derived from this software without specific prior written
    13  * permission.
    14  *
    15  * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND ANY EXPRESS OR
    16  * IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND
    17  * FITNESS FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE
    18  * FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING,
    19  * BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS;
    20  * OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT,
    21  * STRICT LIABILITY, OR TOR (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
    22  * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
    23  *
    24  **************************************************************************************************/
    28 #pragma once
    29 
    30 #include <cutlass/fragment.h>
    31 
    32 namespace cutlass {
    33 namespace gemm {
    34 
    36 
    37 template <typename GlobalIterator_>
    38 struct IgemmSwizzle {
    40  typedef GlobalIterator_ GlobalIterator;
    42  typedef typename GlobalIterator::Fragment Fragment;
    44  typedef typename GlobalIterator::FragmentShape FragmentShape;
    45 
    47  typedef Fragment InputFragment;
    49  typedef Fragment OutputFragment;
    50 
    52  static_assert((platform::is_same<typename Fragment::Element, int8_t>::value), "Works on int8");
    53 
    55  static_assert(FragmentShape::kH % 4 == 0 && ShapeCount<FragmentShape>::kWc % 4 == 0,
    56  "Not multiple of 4");
    57 
    59  CUTLASS_DEVICE IgemmSwizzle() {}
    60 
    62  CUTLASS_DEVICE void transform(Fragment const& src, Fragment& dst) {
    63  // Expose src/dst as int arrays.
    64  int const* src_int = reinterpret_cast<int const*>(&src[0]);
    65  int* dst_int = reinterpret_cast<int*>(&dst[0]);
    66 
    67  // Transpose the data.
    68  for (int d = 0; d < FragmentShape::kD; ++d) {
    69  for (int h = 0; h < FragmentShape::kH / 4; ++h) {
    70  for (int w = 0; w < ShapeCount<FragmentShape>::kWc / 4; ++w) {
    71  int const i0 = d * (ShapeCount<FragmentShape>::kHwc / 4) +
    72  (4 * h + 0) * (ShapeCount<FragmentShape>::kWc / 4) + w;
    73  int const i1 = d * (ShapeCount<FragmentShape>::kHwc / 4) +
    74  (4 * h + 1) * (ShapeCount<FragmentShape>::kWc / 4) + w;
    75  int const i2 = d * (ShapeCount<FragmentShape>::kHwc / 4) +
    76  (4 * h + 2) * (ShapeCount<FragmentShape>::kWc / 4) + w;
    77  int const i3 = d * (ShapeCount<FragmentShape>::kHwc / 4) +
    78  (4 * h + 3) * (ShapeCount<FragmentShape>::kWc / 4) + w;
    79 
    80  int a0 = src_int[i0];
    81  int a1 = src_int[i1];
    82  int a2 = src_int[i2];
    83  int a3 = src_int[i3];
    84 
    85  int b0, b1, b2, b3, c0;
    86  asm volatile("prmt.b32 %0, %1, %2, 0x0040;" : "=r"(b0) : "r"(a0), "r"(a1));
    87  asm volatile("prmt.b32 %0, %1, %2, 0x0040;" : "=r"(c0) : "r"(a2), "r"(a3));
    88  asm volatile("prmt.b32 %0, %1, %2, 0x5410;" : "=r"(b0) : "r"(b0), "r"(c0));
    89 
    90  asm volatile("prmt.b32 %0, %1, %2, 0x0051;" : "=r"(b1) : "r"(a0), "r"(a1));
    91  asm volatile("prmt.b32 %0, %1, %2, 0x0051;" : "=r"(c0) : "r"(a2), "r"(a3));
    92  asm volatile("prmt.b32 %0, %1, %2, 0x5410;" : "=r"(b1) : "r"(b1), "r"(c0));
    93 
    94  asm volatile("prmt.b32 %0, %1, %2, 0x0062;" : "=r"(b2) : "r"(a0), "r"(a1));
    95  asm volatile("prmt.b32 %0, %1, %2, 0x0062;" : "=r"(c0) : "r"(a2), "r"(a3));
    96  asm volatile("prmt.b32 %0, %1, %2, 0x5410;" : "=r"(b2) : "r"(b2), "r"(c0));
    97 
    98  asm volatile("prmt.b32 %0, %1, %2, 0x0073;" : "=r"(b3) : "r"(a0), "r"(a1));
    99  asm volatile("prmt.b32 %0, %1, %2, 0x0073;" : "=r"(c0) : "r"(a2), "r"(a3));
    100  asm volatile("prmt.b32 %0, %1, %2, 0x5410;" : "=r"(b3) : "r"(b3), "r"(c0));
    101 
    102  dst_int[i0] = b0;
    103  dst_int[i1] = b1;
    104  dst_int[i2] = b2;
    105  dst_int[i3] = b3;
    106  }
    107  }
    108  }
    109  }
    110 };
    111 
    113 
    114 } // namespace gemm
    115 } // namespace cutlass
    cutlass
    Definition: convert.h:33
    -
    cutlass::platform::is_same
    std::is_same (false specialization)
    Definition: platform.h:412
    +Go to the documentation of this file.
    1 /***************************************************************************************************
    2  * Copyright (c) 2017-2018, NVIDIA CORPORATION. All rights reserved.
    3  *
    4  * Redistribution and use in source and binary forms, with or without modification, are permitted
    5  * provided that the following conditions are met:
    6  * * Redistributions of source code must retain the above copyright notice, this list of
    7  * conditions and the following disclaimer.
    8  * * Redistributions in binary form must reproduce the above copyright notice, this list of
    9  * conditions and the following disclaimer in the documentation and/or other materials
    10  * provided with the distribution.
    11  * * Neither the name of the NVIDIA CORPORATION nor the names of its contributors may be used
    12  * to endorse or promote products derived from this software without specific prior written
    13  * permission.
    14  *
    15  * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND ANY EXPRESS OR
    16  * IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND
    17  * FITNESS FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE
    18  * FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING,
    19  * BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS;
    20  * OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT,
    21  * STRICT LIABILITY, OR TOR (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
    22  * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
    23  *
    24  **************************************************************************************************/
    28 #pragma once
    29 
    30 #include "cutlass/fragment.h"
    31 
    32 namespace cutlass {
    33 namespace gemm {
    34 
    36 
    37 template <typename GlobalIterator_>
    38 struct IgemmSwizzle {
    40  typedef GlobalIterator_ GlobalIterator;
    42  typedef typename GlobalIterator::Fragment Fragment;
    44  typedef typename GlobalIterator::FragmentShape FragmentShape;
    45 
    47  typedef Fragment InputFragment;
    49  typedef Fragment OutputFragment;
    50 
    52  static_assert((platform::is_same<typename Fragment::Element, int8_t>::value), "Works on int8");
    53 
    55  static_assert(FragmentShape::kH % 4 == 0 && ShapeCount<FragmentShape>::kWc % 4 == 0,
    56  "Not multiple of 4");
    57 
    59  CUTLASS_DEVICE IgemmSwizzle() {}
    60 
    62  CUTLASS_DEVICE void transform(Fragment const& src, Fragment& dst) {
    63  // Expose src/dst as int arrays.
    64  int const* src_int = reinterpret_cast<int const*>(&src[0]);
    65  int* dst_int = reinterpret_cast<int*>(&dst[0]);
    66 
    67  // Transpose the data.
    68  for (int d = 0; d < FragmentShape::kD; ++d) {
    69  for (int h = 0; h < FragmentShape::kH / 4; ++h) {
    70  for (int w = 0; w < ShapeCount<FragmentShape>::kWc / 4; ++w) {
    71  int const i0 = d * (ShapeCount<FragmentShape>::kHwc / 4) +
    72  (4 * h + 0) * (ShapeCount<FragmentShape>::kWc / 4) + w;
    73  int const i1 = d * (ShapeCount<FragmentShape>::kHwc / 4) +
    74  (4 * h + 1) * (ShapeCount<FragmentShape>::kWc / 4) + w;
    75  int const i2 = d * (ShapeCount<FragmentShape>::kHwc / 4) +
    76  (4 * h + 2) * (ShapeCount<FragmentShape>::kWc / 4) + w;
    77  int const i3 = d * (ShapeCount<FragmentShape>::kHwc / 4) +
    78  (4 * h + 3) * (ShapeCount<FragmentShape>::kWc / 4) + w;
    79 
    80  int a0 = src_int[i0];
    81  int a1 = src_int[i1];
    82  int a2 = src_int[i2];
    83  int a3 = src_int[i3];
    84 
    85  // // DEBUG.
    86  // if (threadIdx.x == 0) {
    87  // printf("a=0x%08x 0x%08x 0x%08x 0x%08x\n", a0, a1, a2, a3);
    88  // }
    89 
    90  int b0, b1, b2, b3, c0;
    91  asm volatile("prmt.b32 %0, %1, %2, 0x0040;" : "=r"(b0) : "r"(a0), "r"(a1));
    92  asm volatile("prmt.b32 %0, %1, %2, 0x0040;" : "=r"(c0) : "r"(a2), "r"(a3));
    93  asm volatile("prmt.b32 %0, %1, %2, 0x5410;" : "=r"(b0) : "r"(b0), "r"(c0));
    94 
    95  asm volatile("prmt.b32 %0, %1, %2, 0x0051;" : "=r"(b1) : "r"(a0), "r"(a1));
    96  asm volatile("prmt.b32 %0, %1, %2, 0x0051;" : "=r"(c0) : "r"(a2), "r"(a3));
    97  asm volatile("prmt.b32 %0, %1, %2, 0x5410;" : "=r"(b1) : "r"(b1), "r"(c0));
    98 
    99  asm volatile("prmt.b32 %0, %1, %2, 0x0062;" : "=r"(b2) : "r"(a0), "r"(a1));
    100  asm volatile("prmt.b32 %0, %1, %2, 0x0062;" : "=r"(c0) : "r"(a2), "r"(a3));
    101  asm volatile("prmt.b32 %0, %1, %2, 0x5410;" : "=r"(b2) : "r"(b2), "r"(c0));
    102 
    103  asm volatile("prmt.b32 %0, %1, %2, 0x0073;" : "=r"(b3) : "r"(a0), "r"(a1));
    104  asm volatile("prmt.b32 %0, %1, %2, 0x0073;" : "=r"(c0) : "r"(a2), "r"(a3));
    105  asm volatile("prmt.b32 %0, %1, %2, 0x5410;" : "=r"(b3) : "r"(b3), "r"(c0));
    106 
    107  // // DEBUG.
    108  // if (threadIdx.x == 0) {
    109  // printf("b=0x%08x 0x%08x 0x%08x 0x%08x\n", b0, b1, b2, b3);
    110  // }
    111 
    112  dst_int[i0] = b0;
    113  dst_int[i1] = b1;
    114  dst_int[i2] = b2;
    115  dst_int[i3] = b3;
    116  }
    117  }
    118  }
    119  }
    120 };
    121 
    123 
    124 } // namespace gemm
    125 } // namespace cutlass
    cutlass
    Definition: convert.h:33
    +
    cutlass::platform::is_same
    std::is_same (false specialization)
    Definition: platform.h:420
    cutlass::gemm::IgemmSwizzle::FragmentShape
    GlobalIterator::FragmentShape FragmentShape
    The shape of the source fragment.
    Definition: igemm_swizzle.h:44
    cutlass::gemm::IgemmSwizzle
    Definition: igemm_swizzle.h:38
    cutlass::gemm::IgemmSwizzle::GlobalIterator
    GlobalIterator_ GlobalIterator
    The global iterator.
    Definition: igemm_swizzle.h:40
    cutlass::gemm::IgemmSwizzle::transform
    CUTLASS_DEVICE void transform(Fragment const &src, Fragment &dst)
    Transform a fragment.
    Definition: igemm_swizzle.h:62
    cutlass::gemm::IgemmSwizzle::OutputFragment
    Fragment OutputFragment
    The destination fragment.
    Definition: igemm_swizzle.h:49
    -
    static_assert
    #define static_assert(__e, __m)
    Definition: platform.h:145
    +
    static_assert
    #define static_assert(__e, __m)
    Definition: platform.h:153
    cutlass::gemm::IgemmSwizzle::InputFragment
    Fragment InputFragment
    The source fragment.
    Definition: igemm_swizzle.h:47
    cutlass::gemm::IgemmSwizzle::Fragment
    GlobalIterator::Fragment Fragment
    The source fragment.
    Definition: igemm_swizzle.h:42
    cutlass::gemm::IgemmSwizzle::IgemmSwizzle
    CUTLASS_DEVICE IgemmSwizzle()
    The src/dst must be int8 fragments.
    Definition: igemm_swizzle.h:59
    @@ -92,7 +92,7 @@ $(function() {
    -Generated on Thu May 3 2018 16:36:12 for Cutlass by   +Generated on Wed Sep 19 2018 13:58:50 for Cutlass by   doxygen 1.8.14
    diff --git a/docs/igemm__traits_8h.html b/docs/igemm__traits_8h.html index 32d14d87..897687ee 100644 --- a/docs/igemm__traits_8h.html +++ b/docs/igemm__traits_8h.html @@ -82,34 +82,38 @@ $(function() {

    Defies structural properties of mixed-precision integer GEMM. Multiplicands are assumed to be packed 8bit integers, accumulators are assumed to be 32b signed integers, and output formats vary. More...

    -
    #include <cutlass/convert.h>
    -#include <cutlass/gemm/gemm.h>
    -#include <cutlass/gemm/gemm_epilogue.h>
    -#include <cutlass/gemm/gemm_epilogue_traits.h>
    -#include <cutlass/gemm/gemm_global_tile.h>
    -#include <cutlass/gemm/gemm_shared_tile.h>
    -#include <cutlass/gemm/gemm_traits.h>
    -#include <cutlass/gemm/igemm_epilogue.h>
    -#include <cutlass/gemm/igemm_global_tile.h>
    -#include <cutlass/gemm/igemm_multiply_add.h>
    -#include <cutlass/gemm/igemm_swizzle.h>
    -#include <cutlass/reshape_tile.h>
    +
    #include "cutlass/convert.h"
    +#include "cutlass/gemm/gemm.h"
    +#include "cutlass/gemm/gemm_epilogue.h"
    +#include "cutlass/gemm/gemm_epilogue_traits.h"
    +#include "cutlass/gemm/gemm_global_tile.h"
    +#include "cutlass/gemm/gemm_shared_tile.h"
    +#include "cutlass/gemm/gemm_traits.h"
    +#include "cutlass/gemm/igemm_epilogue.h"
    +#include "cutlass/gemm/igemm_global_tile.h"
    +#include "cutlass/gemm/igemm_multiply_add.h"
    +#include "cutlass/gemm/igemm_swizzle.h"
    +#include "cutlass/reshape_tile.h"

    Go to the source code of this file.

    - + - + - + - + - + - + + + + + @@ -123,13 +127,13 @@ Classes - + - +

    Classes

    struct  cutlass::gemm::IgemmConfig< OutputTile_, ScalarD_, AccumulatorsPerThread_ >
    struct  cutlass::gemm::IgemmConfig< OutputTile_, ScalarD_, ThreadGemmShape_ >
     
    struct  cutlass::gemm::IgemmConfig< OutputTile_, int8_t, AccumulatorsPerThread_ >
    struct  cutlass::gemm::IgemmConfig< OutputTile_, int8_t, ThreadGemmShape_ >
     
    struct  cutlass::gemm::IgemmTileTraitsHelperA< kLayout_, GemmConfig_ >
    struct  cutlass::gemm::IgemmTileTraitsHelperA< kLayout_, GemmConfig_, Index_ >
     
    struct  cutlass::gemm::IgemmTileTraitsHelperA< MatrixLayout::kColumnMajor, GemmConfig_ >
    struct  cutlass::gemm::IgemmTileTraitsHelperA< MatrixLayout::kColumnMajor, GemmConfig_, Index_ >
     
    struct  cutlass::gemm::IgemmTileTraitsHelperB< kLayout_, GemmConfig_ >
    struct  cutlass::gemm::IgemmTileTraitsHelperA< MatrixLayout::kRowMajor, GemmConfig_, Index_ >
     
    struct  cutlass::gemm::IgemmTileTraitsHelperB< MatrixLayout::kRowMajor, GemmConfig_ >
    struct  cutlass::gemm::IgemmTileTraitsHelperB< kLayout_, GemmConfig_, Index_ >
     
    struct  cutlass::gemm::IgemmTileTraitsHelperB< MatrixLayout::kColumnMajor, GemmConfig_, Index_ >
     
    struct  cutlass::gemm::IgemmTileTraitsHelperB< MatrixLayout::kRowMajor, GemmConfig_, Index_ >
     
    struct  cutlass::gemm::IgemmTransformerA< kLayout_, Iterator_ >
     
     
    struct  cutlass::gemm::IgemmTransformerB< MatrixLayout::kRowMajor, Iterator_ >
     
    struct  cutlass::gemm::IgemmTraitsHelper< kLayoutA_, kLayoutB_, OutputTile_, ScalarD_, EpilogueFunctor_, AccumulatorsPerThread_, Index_ >
    struct  cutlass::gemm::IgemmTraitsHelper< kLayoutA_, kLayoutB_, OutputTile_, ScalarD_, EpilogueFunctor_, ThreadGemmShape_, Index_ >
     
    struct  cutlass::gemm::IgemmEpilogueScalar< ScalarD_ >
     
    struct  cutlass::gemm::IgemmEpilogueScalar< int >
     
    struct  cutlass::gemm::IgemmTraits< kLayoutA_, kLayoutB_, OutputTile_, ScalarD_, EpilogueFunctor_, AccumulatorsPerThread_, Index_, Helper_ >
    struct  cutlass::gemm::IgemmTraits< kLayoutA_, kLayoutB_, OutputTile_, ScalarD_, EpilogueFunctor_, ThreadGemmShape_, Index_, Helper_ >
     
    -Generated on Thu May 3 2018 16:36:12 for Cutlass by   +Generated on Wed Sep 19 2018 13:58:50 for Cutlass by   doxygen 1.8.14
    diff --git a/docs/igemm__traits_8h_source.html b/docs/igemm__traits_8h_source.html index ecdd4f1d..e1fa87e4 100644 --- a/docs/igemm__traits_8h_source.html +++ b/docs/igemm__traits_8h_source.html @@ -76,89 +76,108 @@ $(function() {
    igemm_traits.h
    -Go to the documentation of this file.
    1 /***************************************************************************************************
    2  * Copyright (c) 2017-2018, NVIDIA CORPORATION. All rights reserved.
    3  *
    4  * Redistribution and use in source and binary forms, with or without modification, are permitted
    5  * provided that the following conditions are met:
    6  * * Redistributions of source code must retain the above copyright notice, this list of
    7  * conditions and the following disclaimer.
    8  * * Redistributions in binary form must reproduce the above copyright notice, this list of
    9  * conditions and the following disclaimer in the documentation and/or other materials
    10  * provided with the distribution.
    11  * * Neither the name of the NVIDIA CORPORATION nor the names of its contributors may be used
    12  * to endorse or promote products derived from this software without specific prior written
    13  * permission.
    14  *
    15  * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND ANY EXPRESS OR
    16  * IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND
    17  * FITNESS FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE
    18  * FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING,
    19  * BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS;
    20  * OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT,
    21  * STRICT LIABILITY, OR TOR (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
    22  * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
    23  *
    24  **************************************************************************************************/
    30 #pragma once
    31 
    32 #include <cutlass/convert.h>
    33 #include <cutlass/gemm/gemm.h>
    34 #include <cutlass/gemm/gemm_epilogue.h>
    35 #include <cutlass/gemm/gemm_epilogue_traits.h>
    36 #include <cutlass/gemm/gemm_global_tile.h>
    37 #include <cutlass/gemm/gemm_shared_tile.h>
    38 #include <cutlass/gemm/gemm_traits.h>
    39 #include <cutlass/gemm/igemm_epilogue.h>
    40 #include <cutlass/gemm/igemm_global_tile.h>
    41 #include <cutlass/gemm/igemm_multiply_add.h>
    42 #include <cutlass/gemm/igemm_swizzle.h>
    43 #include <cutlass/reshape_tile.h>
    44 
    45 namespace cutlass {
    46 namespace gemm {
    47 
    49 
    50 template <
    52  typename OutputTile_,
    54  typename ScalarD_,
    56  typename AccumulatorsPerThread_>
    57 struct IgemmConfig
    58  : public GemmConfig<
    60  int8_t,
    62  int8_t,
    64  ScalarD_,
    66  ScalarD_,
    68  OutputTile_,
    70  ThreadMultiplyAdd<AccumulatorsPerThread_, Shape<1, 4, 8>, int8_t, int8_t, int>,
    72  4,
    74  4,
    76  16,
    78  4,
    80  4,
    82  16,
    84  1,
    86  4,
    88  1,
    90  2> {};
    91 
    93 
    94 template <typename OutputTile_, typename AccumulatorsPerThread_>
    95 struct IgemmConfig<OutputTile_, int8_t, AccumulatorsPerThread_>
    96  : public GemmConfig<
    98  int8_t,
    100  int8_t,
    102  int8_t,
    104  int8_t,
    106  OutputTile_,
    108  ThreadMultiplyAdd<AccumulatorsPerThread_, Shape<1, 4, 8>, int8_t, int8_t, int>,
    110  4,
    112  4,
    114  16,
    116  4,
    118  4,
    120  16,
    122  4,
    124  4,
    126  4,
    128  2> {};
    129 
    131 
    132 template <enum MatrixLayout::Kind kLayout_, typename GemmConfig_>
    133 struct IgemmTileTraitsHelperA : public GemmTileTraitsHelperA<kLayout_, GemmConfig_> {};
    134 
    136 
    137 template <typename GemmConfig_>
    138 struct IgemmTileTraitsHelperA<MatrixLayout::kColumnMajor, GemmConfig_>
    139  : public GemmTileTraitsHelperA<MatrixLayout::kColumnMajor, GemmConfig_> {
    141  typedef GemmTileTraitsHelperA<MatrixLayout::kColumnMajor, GemmConfig_> Base;
    142 
    144  static int const kScalarsPerStsA = 16;
    145 
    147  typedef IgemmContiguousGlobalTileTraits<
    148  GemmOperand::kA,
    149  // The layout.
    150  MatrixLayout::kColumnMajor,
    151  // The pointer is float const.
    152  int8_t const,
    153  // The tile has size KxM in GEMM's terminology.
    154  Shape<1, GemmConfig_::OutputTile::kD, GemmConfig_::OutputTile::kW>,
    155  // The threads are distributed as warps x 32 (the traits may reorganize).
    156  Shape<1, ShapeCount<typename GemmConfig_::Warps>::kCount, GemmConfig_::kWarpSize>,
    157  // The number of scalars per LDG (LDG.32 or LDG.128, etc).
    158  4>
    159  GlobalTileTraits;
    160 
    162  typedef GemmSharedStoreTileAbTraits<
    163  // The pointer is float.
    164  int8_t,
    165  // The tile has size KxM in GEMM's terminology.
    166  Shape<GemmConfig_::kStages, GemmConfig_::OutputTile::kD / 4, GemmConfig_::OutputTile::kW * 4>,
    167  // The threads are distributed as warps x 32 (the traits may reorganize).
    168  typename GlobalTileTraits::Threads,
    169  // The number of scalars per STS (STS.32 or STS.128, etc).
    170  kScalarsPerStsA>
    171  SharedStoreTileTraits;
    172 };
    173 
    175 
    176 template <enum MatrixLayout::Kind kLayout_, typename GemmConfig_>
    177 struct IgemmTileTraitsHelperB : public GemmTileTraitsHelperB<kLayout_, GemmConfig_> {};
    178 
    180 
    181 template <typename GemmConfig_>
    182 struct IgemmTileTraitsHelperB<MatrixLayout::kRowMajor, GemmConfig_>
    183  : public GemmTileTraitsHelperB<MatrixLayout::kRowMajor, GemmConfig_> {
    185  typedef GemmTileTraitsHelperB<MatrixLayout::kRowMajor, GemmConfig_> Base;
    186 
    188  static int const kScalarsPerStsB = 16;
    189 
    191  typedef IgemmContiguousGlobalTileTraits<
    192  GemmOperand::kB,
    193  // The layout.
    194  MatrixLayout::kRowMajor,
    195  // The pointer is float const.
    196  int8_t const,
    197  // The tile has size KxM in GEMM's terminology.
    198  Shape<1, GemmConfig_::OutputTile::kD, GemmConfig_::OutputTile::kH>,
    199  // The threads are distributed as warps x 32 (the traits may reorganize).
    200  Shape<1, ShapeCount<typename GemmConfig_::Warps>::kCount, GemmConfig_::kWarpSize>,
    201  // The number of scalars per LDG (LDG.32 or LDG.128, etc).
    202  4>
    203  GlobalTileTraits;
    204 
    206  typedef GemmSharedStoreTileAbTraits<
    207  // The pointer is float.
    208  int8_t,
    209  // The tile has size KxM in GEMM's terminology.
    210  Shape<GemmConfig_::kStages, GemmConfig_::OutputTile::kD / 4, GemmConfig_::OutputTile::kH * 4>,
    211  // The threads are distributed as warps x 32 (the traits may reorganize).
    212  typename GlobalTileTraits::Threads,
    213  // The number of scalars per STS (STS.32 or STS.128, etc).
    214  kScalarsPerStsB>
    215  SharedStoreTileTraits;
    216 };
    217 
    219 
    220 template <enum MatrixLayout::Kind kLayout_, typename Iterator_>
    221 struct IgemmTransformerA {};
    222 
    223 template <typename Iterator_>
    224 struct IgemmTransformerA<MatrixLayout::kRowMajor, Iterator_> {
    225  typedef Copy<typename Iterator_::Fragment> Transformer;
    226 };
    227 
    228 template <typename Iterator_>
    229 struct IgemmTransformerA<MatrixLayout::kColumnMajor, Iterator_> {
    230  typedef IgemmSwizzle<Iterator_> Transformer;
    231 };
    232 
    234 
    235 template <enum MatrixLayout::Kind kLayout_, typename Iterator_>
    236 struct IgemmTransformerB {};
    237 
    238 template <typename Iterator_>
    239 struct IgemmTransformerB<MatrixLayout::kColumnMajor, Iterator_> {
    240  typedef Copy<typename Iterator_::Fragment> Transformer;
    241 };
    242 
    243 template <typename Iterator_>
    244 struct IgemmTransformerB<MatrixLayout::kRowMajor, Iterator_> {
    245  typedef IgemmSwizzle<Iterator_> Transformer;
    246 };
    247 
    249 
    250 template <
    252  MatrixLayout::Kind kLayoutA_,
    254  MatrixLayout::Kind kLayoutB_,
    256  typename OutputTile_,
    258  typename ScalarD_,
    260  typename EpilogueFunctor_,
    262  typename AccumulatorsPerThread_ = Shape<32, 8, 8>,
    264  typename Index_ = int>
    265 struct IgemmTraitsHelper {
    267  typedef IgemmConfig<OutputTile_, ScalarD_, AccumulatorsPerThread_> GemmConfig;
    269  typedef IgemmTileTraitsHelperA<kLayoutA_, GemmConfig> GemmTileTraitsHelperA;
    271  typedef IgemmTileTraitsHelperB<kLayoutB_, GemmConfig> GemmTileTraitsHelperB;
    272 
    274  typedef GemmGlobalIteratorAb<typename GemmTileTraitsHelperA::GlobalTileTraits, Index_>
    275  GlobalLoadIteratorA;
    277  typedef typename IgemmTransformerA<GemmTileTraitsHelperA::kLayout,
    278  GlobalLoadIteratorA>::Transformer GlobalTransformerA;
    280  typedef TileStoreIterator<typename GemmTileTraitsHelperA::SharedStoreTileTraits,
    281  typename GemmTileTraitsHelperA::SharedStoreTileTraits::Scalar,
    282  IteratorAdvance::kH,
    283  MemorySpace::kShared>
    284  SharedStoreIteratorA;
    286  typedef GlobalLoadStream<GlobalLoadIteratorA, SharedStoreIteratorA, GlobalTransformerA>
    287  GlobalLoadStreamA;
    288 
    290  typedef GemmGlobalIteratorAb<typename GemmTileTraitsHelperB::GlobalTileTraits, Index_>
    291  GlobalLoadIteratorB;
    292  // The default transformer for B.
    293  typedef typename IgemmTransformerB<GemmTileTraitsHelperB::kLayout,
    294  GlobalLoadIteratorB>::Transformer GlobalTransformerB;
    296  typedef TileStoreIterator<typename GemmTileTraitsHelperB::SharedStoreTileTraits,
    297  typename GemmTileTraitsHelperB::SharedStoreTileTraits::Scalar,
    298  IteratorAdvance::kH,
    299  MemorySpace::kShared>
    300  SharedStoreIteratorB;
    302  typedef GlobalLoadStream<GlobalLoadIteratorB, SharedStoreIteratorB, GlobalTransformerB>
    303  GlobalLoadStreamB;
    304 
    306  typedef TileLoadIterator<typename GemmTileTraitsHelperA::SharedLoadTileTraits,
    307  typename GemmTileTraitsHelperA::SharedLoadTileTraits::Scalar,
    308  IteratorAdvance::kH,
    309  MemorySpace::kShared>
    310  SharedLoadIteratorA;
    312  typedef SharedLoadStream<SharedLoadIteratorA, Copy<typename SharedLoadIteratorA::Fragment> >
    313  SharedLoadStreamA;
    315  typedef TileLoadIterator<typename GemmTileTraitsHelperB::SharedLoadTileTraits,
    316  typename GemmTileTraitsHelperB::SharedLoadTileTraits::Scalar,
    317  IteratorAdvance::kH,
    318  MemorySpace::kShared>
    319  SharedLoadIteratorB;
    321  typedef SharedLoadStream<SharedLoadIteratorB, Copy<typename SharedLoadIteratorB::Fragment> >
    322  SharedLoadStreamB;
    323 
    325  typedef typename GemmConfig::MultiplyAdd MultiplyAdd;
    327  typedef ClearAccumulators<typename MultiplyAdd::ScalarC> ClearAccumulators;
    328 
    330  typedef IgemmEpilogue<IgemmEpilogueTraits<GemmConfig, EpilogueFunctor_> > Epilogue;
    331 };
    332 
    334 
    335 template <typename ScalarD_>
    336 struct IgemmEpilogueScalar {
    337  typedef float Scalar;
    338 };
    339 
    340 template <>
    341 struct IgemmEpilogueScalar<int> {
    342  typedef int Scalar;
    343 };
    344 
    346 
    347 template <
    349  MatrixLayout::Kind kLayoutA_,
    351  MatrixLayout::Kind kLayoutB_,
    353  typename OutputTile_ = Shape<32, 128, 128>,
    355  typename ScalarD_ = int,
    357  typename EpilogueFunctor_ = LinearScaling<typename IgemmEpilogueScalar<ScalarD_>::Scalar>,
    359  typename AccumulatorsPerThread_ = Shape<32, 8, 8>,
    361  typename Index_ = int,
    363  typename Helper_ = IgemmTraitsHelper<kLayoutA_,
    364  kLayoutB_,
    365  OutputTile_,
    366  ScalarD_,
    367  EpilogueFunctor_,
    368  AccumulatorsPerThread_,
    369  Index_> >
    370 struct IgemmTraits : public GemmTraits<
    371  // The config.
    372  typename Helper_::GemmConfig,
    373  // The stream to load A from global memory to shared memory.
    374  typename Helper_::GlobalLoadStreamA,
    375  // The stream to load B from global memory to shared memory.
    376  typename Helper_::GlobalLoadStreamB,
    377  // The stream to load A from shared memory.
    378  typename Helper_::SharedLoadStreamA,
    379  // The stream to load B from shared memory.
    380  typename Helper_::SharedLoadStreamB,
    381  // The epilogue.
    382  typename Helper_::Epilogue,
    383  // The block swizzle to reorganize the grid.
    384  IdentityBlockSwizzle,
    385  // The index.
    386  Index_,
    387  // The tool used to clear accumulators.
    388  typename Helper_::ClearAccumulators> {};
    389 
    391 
    392 } // namespace gemm
    393 } // namespace cutlass
    cutlass::MemorySpace::kShared
    Definition: load_store.h:42
    -
    cutlass::gemm::IgemmTraitsHelper::SharedLoadIteratorB
    TileLoadIterator< typename GemmTileTraitsHelperB::SharedLoadTileTraits, typename GemmTileTraitsHelperB::SharedLoadTileTraits::Scalar, IteratorAdvance::kH, MemorySpace::kShared > SharedLoadIteratorB
    The iterator to load B from shared memory.
    Definition: igemm_traits.h:319
    +Go to the documentation of this file.
    1 /***************************************************************************************************
    2  * Copyright (c) 2017-2018, NVIDIA CORPORATION. All rights reserved.
    3  *
    4  * Redistribution and use in source and binary forms, with or without modification, are permitted
    5  * provided that the following conditions are met:
    6  * * Redistributions of source code must retain the above copyright notice, this list of
    7  * conditions and the following disclaimer.
    8  * * Redistributions in binary form must reproduce the above copyright notice, this list of
    9  * conditions and the following disclaimer in the documentation and/or other materials
    10  * provided with the distribution.
    11  * * Neither the name of the NVIDIA CORPORATION nor the names of its contributors may be used
    12  * to endorse or promote products derived from this software without specific prior written
    13  * permission.
    14  *
    15  * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND ANY EXPRESS OR
    16  * IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND
    17  * FITNESS FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE
    18  * FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING,
    19  * BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS;
    20  * OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT,
    21  * STRICT LIABILITY, OR TOR (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
    22  * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
    23  *
    24  **************************************************************************************************/
    30 #pragma once
    31 
    32 #include "cutlass/convert.h"
    33 #include "cutlass/gemm/gemm.h"
    34 #include "cutlass/gemm/gemm_epilogue.h"
    35 #include "cutlass/gemm/gemm_epilogue_traits.h"
    36 #include "cutlass/gemm/gemm_global_tile.h"
    37 #include "cutlass/gemm/gemm_shared_tile.h"
    38 #include "cutlass/gemm/gemm_traits.h"
    39 #include "cutlass/gemm/igemm_epilogue.h"
    40 #include "cutlass/gemm/igemm_global_tile.h"
    41 #include "cutlass/gemm/igemm_multiply_add.h"
    42 #include "cutlass/gemm/igemm_swizzle.h"
    43 #include "cutlass/reshape_tile.h"
    44 
    45 namespace cutlass {
    46 namespace gemm {
    47 
    49 
    50 template <
    52  typename OutputTile_,
    54  typename ScalarD_,
    56  typename ThreadGemmShape_>
    57 struct IgemmConfig : public GemmConfig<
    59  int8_t,
    61  int8_t,
    63  ScalarD_,
    65  ScalarD_,
    67  OutputTile_,
    69  ThreadMultiplyAdd<ThreadGemmShape_, Shape<1, 4, 8>, int8_t, int8_t, int>,
    71  4,
    73  4,
    75  16,
    77  4,
    79  4,
    81  16,
    83  1,
    85  4,
    87  1,
    89  2,
    91  false,
    93  false,
    95  false> {};
    96 
    98 
    99 template <typename OutputTile_, typename ThreadGemmShape_>
    100 struct IgemmConfig<OutputTile_, int8_t, ThreadGemmShape_>
    101  : public GemmConfig<
    103  int8_t,
    105  int8_t,
    107  int8_t,
    109  int8_t,
    111  OutputTile_,
    113  ThreadMultiplyAdd<ThreadGemmShape_, Shape<1, 4, 8>, int8_t, int8_t, int>,
    115  4,
    117  4,
    119  16,
    121  4,
    123  4,
    125  16,
    127  4,
    129  4,
    131  4,
    133  2,
    135  false,
    137  true,
    139  false> {};
    140 
    142 
    143 template <enum MatrixLayout::Kind kLayout_, typename GemmConfig_, typename Index_>
    144 struct IgemmTileTraitsHelperA : public GemmTileTraitsHelperA<kLayout_, GemmConfig_> {};
    145 
    147 
    148 template <typename GemmConfig_, typename Index_>
    149 struct IgemmTileTraitsHelperA<MatrixLayout::kColumnMajor, GemmConfig_, Index_>
    150  : public GemmTileTraitsHelperA<MatrixLayout::kColumnMajor, GemmConfig_> {
    152  typedef GemmTileTraitsHelperA<MatrixLayout::kColumnMajor, GemmConfig_> Base;
    153 
    155  static int const kScalarsPerStsA = 16;
    156 
    158  typedef IgemmGlobalTileTraits<
    159  GemmOperand::kA,
    160  // The layout.
    161  MatrixLayout::kColumnMajor,
    162  // The pointer is float const.
    163  int8_t const,
    164  // The tile has size KxM in GEMM's terminology.
    165  Shape<1, GemmConfig_::OutputTile::kD, GemmConfig_::OutputTile::kW>,
    166  // The threads are distributed as warps x 32 (the traits may reorganize).
    167  Shape<1, ShapeCount<typename GemmConfig_::Warps>::kCount, GemmConfig_::kWarpSize>,
    168  // The number of scalars per LDG (LDG.32 or LDG.128, etc).
    169  GemmConfig_::kScalarsPerLdgA>
    170  GlobalTileTraits;
    171 
    173  typedef GemmGlobalIteratorAb<GlobalTileTraits, Index_> GlobalLoadIterator;
    174 
    176  typedef GemmSharedStoreTileAbTraits<
    177  // The pointer is float.
    178  int8_t,
    179  // The tile has size KxM in GEMM's terminology.
    180  Shape<GemmConfig_::kStages, GemmConfig_::OutputTile::kD / 4, GemmConfig_::OutputTile::kW * 4>,
    181  // The threads are distributed as warps x 32 (the traits may reorganize).
    182  typename GlobalTileTraits::Threads,
    183  // The number of scalars per STS (STS.32 or STS.128, etc).
    184  kScalarsPerStsA>
    185  SharedStoreTileTraits;
    186 };
    187 
    189 
    190 template <typename GemmConfig_, typename Index_>
    191 struct IgemmTileTraitsHelperA<MatrixLayout::kRowMajor, GemmConfig_, Index_> {
    193  static MatrixLayout::Kind const kLayout = MatrixLayout::kRowMajor;
    194 
    196  typedef int8_t Scalar;
    198  typedef int8_t MultiplyAddScalar;
    199 
    201  static int const kScalarsPerStsA = 16;
    202 
    204  typedef IgemmGlobalTileTraits<
    205  GemmOperand::kA,
    206  // The layout.
    207  MatrixLayout::kRowMajor,
    208  // The pointer is float const.
    209  int8_t const,
    210  // The tile has size NxK in GEMM's terminology.
    211  Shape<1, GemmConfig_::OutputTile::kW, GemmConfig_::OutputTile::kD>,
    212  // The threads are distributed as warps x 32 (the traits may reorganize).
    213  Shape<1, ShapeCount<typename GemmConfig_::Warps>::kCount, GemmConfig_::kWarpSize>,
    214  // The number of scalars per LDG (LDG.32 or LDG.128, etc).
    215  GemmConfig_::kScalarsPerLdgA>
    216  GlobalTileTraits;
    217 
    219  typedef IgemmGlobalIteratorAb<GlobalTileTraits, Index_> GlobalLoadIterator;
    220 
    222  typedef GemmSharedStoreWithSkewTileAbTraits<
    223  // The pointer is int8.
    224  int8_t,
    225  // The tile has size KxN in GEMM's terminology.
    226  Shape<GemmConfig_::kStages, GemmConfig_::OutputTile::kD / 4, GemmConfig_::OutputTile::kW * 4>,
    227  // The threads are distributed as (threads / K) x K (the traits may reorganize).
    228  typename GlobalTileTraits::Threads,
    229  // The number of scalars per STS.
    230  kScalarsPerStsA,
    231  // The skew to avoid bank conflicts added in the tile W dimension.
    232  16>
    233  SharedStoreTileTraits;
    234 
    236  typedef GemmSharedLoadTileATraits<
    237  // The pointer is float const.
    238  int8_t const,
    239  // The output tile size.
    240  typename GemmConfig_::OutputTile,
    241  // The number of warps.
    242  typename GemmConfig_::Warps,
    243  // The number of threads per warp.
    244  typename GemmConfig_::MultiplyAdd::ThreadsPerWarp,
    245  // The shape of the FMA instruction.
    246  typename GemmConfig_::InstructionShape,
    247  // The number of stages.
    248  GemmConfig_::kStages,
    249  // The number of scalars per LDS.
    250  16,
    251  // The skew.
    252  SharedStoreTileTraits::kSkew>
    253  SharedLoadTileTraits;
    254 };
    255 
    257 
    258 template <enum MatrixLayout::Kind kLayout_, typename GemmConfig_, typename Index_>
    259 struct IgemmTileTraitsHelperB : public GemmTileTraitsHelperB<kLayout_, GemmConfig_> {};
    260 
    262 
    263 template <typename GemmConfig_, typename Index_>
    264 struct IgemmTileTraitsHelperB<MatrixLayout::kColumnMajor, GemmConfig_, Index_> {
    266  static MatrixLayout::Kind const kLayout = MatrixLayout::kColumnMajor;
    267 
    269  typedef int8_t Scalar;
    271  typedef int8_t MultiplyAddScalar;
    272 
    274  static int const kScalarsPerStsB = 16;
    275 
    277  typedef IgemmGlobalTileTraits<
    278  GemmOperand::kB,
    279  // The layout.
    280  MatrixLayout::kColumnMajor,
    281  // The pointer is float const.
    282  int8_t const,
    283  // The tile has size NxK in GEMM's terminology.
    284  Shape<1, GemmConfig_::OutputTile::kH, GemmConfig_::OutputTile::kD>,
    285  // The threads are distributed as warps x 32 (the traits may reorganize).
    286  Shape<1, ShapeCount<typename GemmConfig_::Warps>::kCount, GemmConfig_::kWarpSize>,
    287  // The number of scalars per LDG (LDG.32 or LDG.128, etc).
    288  GemmConfig_::kScalarsPerLdgB>
    289  GlobalTileTraits;
    290 
    292  typedef IgemmGlobalIteratorAb<GlobalTileTraits, Index_> GlobalLoadIterator;
    293 
    295  typedef GemmSharedStoreWithSkewTileAbTraits<
    296  // The pointer is int8.
    297  int8_t,
    298  // The tile has size KxN in GEMM's terminology.
    299  Shape<GemmConfig_::kStages, GemmConfig_::OutputTile::kD / 4, GemmConfig_::OutputTile::kH * 4>,
    300  // The threads are distributed as (threads / K) x K (the traits may reorganize).
    301  typename GlobalTileTraits::Threads,
    302  // The number of scalars per STS.
    303  kScalarsPerStsB,
    304  // The skew to avoid bank conflicts added in the tile W dimension.
    305  16>
    306  SharedStoreTileTraits;
    307 
    309  typedef GemmSharedLoadTileBTraits<
    310  // The pointer is float const.
    311  int8_t const,
    312  // The output tile size.
    313  typename GemmConfig_::OutputTile,
    314  // The number of warps.
    315  typename GemmConfig_::Warps,
    316  // The number of threads per warp.
    317  typename GemmConfig_::MultiplyAdd::ThreadsPerWarp,
    318  // The shape of the FMA instruction.
    319  typename GemmConfig_::InstructionShape,
    320  // The number of stages.
    321  GemmConfig_::kStages,
    322  // The number of scalars per LDS.
    323  16,
    324  // The skew.
    325  SharedStoreTileTraits::kSkew>
    326  SharedLoadTileTraits;
    327 };
    328 
    330 
    331 template <typename GemmConfig_, typename Index_>
    332 struct IgemmTileTraitsHelperB<MatrixLayout::kRowMajor, GemmConfig_, Index_>
    333  : public GemmTileTraitsHelperB<MatrixLayout::kRowMajor, GemmConfig_> {
    335  typedef GemmTileTraitsHelperB<MatrixLayout::kRowMajor, GemmConfig_> Base;
    336 
    338  static int const kScalarsPerStsB = 16;
    339 
    341  typedef IgemmGlobalTileTraits<
    342  GemmOperand::kB,
    343  // The layout.
    344  MatrixLayout::kRowMajor,
    345  // The pointer is float const.
    346  int8_t const,
    347  // The tile has size KxM in GEMM's terminology.
    348  Shape<1, GemmConfig_::OutputTile::kD, GemmConfig_::OutputTile::kH>,
    349  // The threads are distributed as warps x 32 (the traits may reorganize).
    350  Shape<1, ShapeCount<typename GemmConfig_::Warps>::kCount, GemmConfig_::kWarpSize>,
    351  // The number of scalars per LDG (LDG.32 or LDG.128, etc).
    352  GemmConfig_::kScalarsPerLdgB>
    353  GlobalTileTraits;
    354 
    356  typedef GemmGlobalIteratorAb<GlobalTileTraits, Index_> GlobalLoadIterator;
    357 
    359  typedef GemmSharedStoreTileAbTraits<
    360  // The pointer is float.
    361  int8_t,
    362  // The tile has size KxM in GEMM's terminology.
    363  Shape<GemmConfig_::kStages, GemmConfig_::OutputTile::kD / 4, GemmConfig_::OutputTile::kH * 4>,
    364  // The threads are distributed as warps x 32 (the traits may reorganize).
    365  typename GlobalTileTraits::Threads,
    366  // The number of scalars per STS (STS.32 or STS.128, etc).
    367  kScalarsPerStsB>
    368  SharedStoreTileTraits;
    369 };
    370 
    372 
    373 template <enum MatrixLayout::Kind kLayout_, typename Iterator_>
    374 struct IgemmTransformerA {};
    375 
    376 template <typename Iterator_>
    377 struct IgemmTransformerA<MatrixLayout::kRowMajor, Iterator_> {
    378  typedef Copy<typename Iterator_::Fragment> Transformer;
    379 };
    380 
    381 template <typename Iterator_>
    382 struct IgemmTransformerA<MatrixLayout::kColumnMajor, Iterator_> {
    383  typedef IgemmSwizzle<Iterator_> Transformer;
    384 };
    385 
    387 
    388 template <enum MatrixLayout::Kind kLayout_, typename Iterator_>
    389 struct IgemmTransformerB {};
    390 
    391 template <typename Iterator_>
    392 struct IgemmTransformerB<MatrixLayout::kColumnMajor, Iterator_> {
    393  typedef Copy<typename Iterator_::Fragment> Transformer;
    394 };
    395 
    396 template <typename Iterator_>
    397 struct IgemmTransformerB<MatrixLayout::kRowMajor, Iterator_> {
    398  typedef IgemmSwizzle<Iterator_> Transformer;
    399 };
    400 
    402 
    403 template <
    405  MatrixLayout::Kind kLayoutA_,
    407  MatrixLayout::Kind kLayoutB_,
    409  typename OutputTile_,
    411  typename ScalarD_,
    413  typename EpilogueFunctor_,
    415  typename ThreadGemmShape_ = Shape<32, 8, 8>,
    417  typename Index_ = int>
    418 struct IgemmTraitsHelper {
    420  typedef IgemmConfig<OutputTile_, ScalarD_, ThreadGemmShape_> GemmConfig;
    422  typedef IgemmTileTraitsHelperA<kLayoutA_, GemmConfig, Index_> GemmTileTraitsHelperA;
    424  typedef IgemmTileTraitsHelperB<kLayoutB_, GemmConfig, Index_> GemmTileTraitsHelperB;
    425 
    427  typedef typename GemmTileTraitsHelperA::GlobalLoadIterator GlobalLoadIteratorA;
    429  typedef typename IgemmTransformerA<GemmTileTraitsHelperA::kLayout,
    430  GlobalLoadIteratorA>::Transformer GlobalTransformerA;
    432  typedef TileStoreIterator<typename GemmTileTraitsHelperA::SharedStoreTileTraits,
    433  typename GemmTileTraitsHelperA::SharedStoreTileTraits::Scalar,
    434  IteratorAdvance::kH,
    435  MemorySpace::kShared>
    436  SharedStoreIteratorA;
    438  typedef GlobalLoadStream<GemmOperand::kA,
    439  GlobalLoadIteratorA,
    440  SharedStoreIteratorA,
    441  GlobalTransformerA>
    442  GlobalLoadStreamA;
    443 
    445  typedef typename GemmTileTraitsHelperB::GlobalLoadIterator GlobalLoadIteratorB;
    446  // The default transformer for B.
    447  typedef typename IgemmTransformerB<GemmTileTraitsHelperB::kLayout,
    448  GlobalLoadIteratorB>::Transformer GlobalTransformerB;
    450  typedef TileStoreIterator<typename GemmTileTraitsHelperB::SharedStoreTileTraits,
    451  typename GemmTileTraitsHelperB::SharedStoreTileTraits::Scalar,
    452  IteratorAdvance::kH,
    453  MemorySpace::kShared>
    454  SharedStoreIteratorB;
    456  typedef GlobalLoadStream<GemmOperand::kB,
    457  GlobalLoadIteratorB,
    458  SharedStoreIteratorB,
    459  GlobalTransformerB>
    460  GlobalLoadStreamB;
    461 
    463  typedef TileLoadIterator<typename GemmTileTraitsHelperA::SharedLoadTileTraits,
    464  typename GemmTileTraitsHelperA::SharedLoadTileTraits::Scalar,
    465  IteratorAdvance::kH,
    466  MemorySpace::kShared>
    467  SharedLoadIteratorA;
    469  typedef SharedLoadStream<SharedLoadIteratorA, Copy<typename SharedLoadIteratorA::Fragment> >
    470  SharedLoadStreamA;
    472  typedef TileLoadIterator<typename GemmTileTraitsHelperB::SharedLoadTileTraits,
    473  typename GemmTileTraitsHelperB::SharedLoadTileTraits::Scalar,
    474  IteratorAdvance::kH,
    475  MemorySpace::kShared>
    476  SharedLoadIteratorB;
    478  typedef SharedLoadStream<SharedLoadIteratorB, Copy<typename SharedLoadIteratorB::Fragment> >
    479  SharedLoadStreamB;
    480 
    482  typedef typename GemmConfig::MultiplyAdd MultiplyAdd;
    484  typedef ClearAccumulators<typename MultiplyAdd::ScalarC> ClearAccumulators;
    485 
    487  typedef IgemmEpilogue<IgemmEpilogueTraits<GemmConfig, EpilogueFunctor_> > Epilogue;
    488 };
    489 
    491 
    492 template <typename ScalarD_>
    493 struct IgemmEpilogueScalar {
    494  typedef float Scalar;
    495 };
    496 
    497 template <>
    498 struct IgemmEpilogueScalar<int> {
    499  typedef int Scalar;
    500 };
    501 
    503 
    504 template <
    506  MatrixLayout::Kind kLayoutA_,
    508  MatrixLayout::Kind kLayoutB_,
    510  typename OutputTile_ = Shape<32, 128, 128>,
    512  typename ScalarD_ = int,
    514  typename EpilogueFunctor_ = LinearScaling<typename IgemmEpilogueScalar<ScalarD_>::Scalar>,
    516  typename ThreadGemmShape_ = Shape<32, 8, 8>,
    518  typename Index_ = int,
    520  typename Helper_ = IgemmTraitsHelper<kLayoutA_,
    521  kLayoutB_,
    522  OutputTile_,
    523  ScalarD_,
    524  EpilogueFunctor_,
    525  ThreadGemmShape_,
    526  Index_> >
    527 struct IgemmTraits : public GemmTraits<
    528  // The config.
    529  typename Helper_::GemmConfig,
    530  // The stream to load A from global memory to shared memory.
    531  typename Helper_::GlobalLoadStreamA,
    532  // The stream to load B from global memory to shared memory.
    533  typename Helper_::GlobalLoadStreamB,
    534  // The stream to load A from shared memory.
    535  typename Helper_::SharedLoadStreamA,
    536  // The stream to load B from shared memory.
    537  typename Helper_::SharedLoadStreamB,
    538  // The epilogue.
    539  typename Helper_::Epilogue,
    540  // The block swizzle to reorganize the grid.
    541  IdentityBlockSwizzle,
    542  // The index.
    543  Index_,
    544  // The tool used to clear accumulators.
    545  typename Helper_::ClearAccumulators> {};
    546 
    548 
    549 } // namespace gemm
    550 } // namespace cutlass
    cutlass::gemm::IgemmTraitsHelper::GlobalTransformerB
    IgemmTransformerB< GemmTileTraitsHelperB::kLayout, GlobalLoadIteratorB >::Transformer GlobalTransformerB
    Definition: igemm_traits.h:448
    +
    cutlass::MemorySpace::kShared
    Definition: load_store.h:41
    +
    cutlass::gemm::IgemmTileTraitsHelperB< MatrixLayout::kRowMajor, GemmConfig_, Index_ >::Base
    GemmTileTraitsHelperB< MatrixLayout::kRowMajor, GemmConfig_ > Base
    The base config.
    Definition: igemm_traits.h:335
    cutlass
    Definition: convert.h:33
    -
    cutlass::gemm::IgemmTransformerA< MatrixLayout::kColumnMajor, Iterator_ >::Transformer
    IgemmSwizzle< Iterator_ > Transformer
    Definition: igemm_traits.h:230
    +
    cutlass::gemm::GemmSharedLoadTileATraits
    Definition: gemm_shared_tile.h:128
    +
    cutlass::gemm::IgemmGlobalTileTraits::Threads
    Base::Threads Threads
    The threads.
    Definition: igemm_global_tile.h:66
    +
    cutlass::gemm::IgemmTraitsHelper::GemmTileTraitsHelperB
    IgemmTileTraitsHelperB< kLayoutB_, GemmConfig, Index_ > GemmTileTraitsHelperB
    The GEMM config for B.
    Definition: igemm_traits.h:424
    + +
    cutlass::gemm::IgemmTransformerA< MatrixLayout::kColumnMajor, Iterator_ >::Transformer
    IgemmSwizzle< Iterator_ > Transformer
    Definition: igemm_traits.h:383
    gemm_global_tile.h
    Defines iterators for efficiently loading and storing to global memory.
    -
    cutlass::gemm::IgemmTraitsHelper::GlobalLoadIteratorA
    GemmGlobalIteratorAb< typename GemmTileTraitsHelperA::GlobalTileTraits, Index_ > GlobalLoadIteratorA
    The iterator to load A from global memory.
    Definition: igemm_traits.h:275
    igemm_swizzle.h
    Transposes a fragment of data containing packed 8-bit integer elements.
    -
    cutlass::gemm::IgemmTransformerB< MatrixLayout::kColumnMajor, Iterator_ >::Transformer
    Copy< typename Iterator_::Fragment > Transformer
    Definition: igemm_traits.h:240
    +
    cutlass::gemm::IgemmTransformerB< MatrixLayout::kColumnMajor, Iterator_ >::Transformer
    Copy< typename Iterator_::Fragment > Transformer
    Definition: igemm_traits.h:393
    +
    cutlass::gemm::IgemmTileTraitsHelperA< MatrixLayout::kRowMajor, GemmConfig_, Index_ >::SharedStoreTileTraits
    GemmSharedStoreWithSkewTileAbTraits< int8_t, Shape< GemmConfig_::kStages, GemmConfig_::OutputTile::kD/4, GemmConfig_::OutputTile::kW *4 >, typename GlobalTileTraits::Threads, kScalarsPerStsA, 16 > SharedStoreTileTraits
    The traits class to build the iterator to store data to shared memory for A^N.
    Definition: igemm_traits.h:233
    +
    cutlass::gemm::IgemmTileTraitsHelperB< MatrixLayout::kColumnMajor, GemmConfig_, Index_ >::GlobalTileTraits
    IgemmGlobalTileTraits< GemmOperand::kB, MatrixLayout::kColumnMajor, int8_t const, Shape< 1, GemmConfig_::OutputTile::kH, GemmConfig_::OutputTile::kD >, Shape< 1, ShapeCount< typename GemmConfig_::Warps >::kCount, GemmConfig_::kWarpSize >, GemmConfig_::kScalarsPerLdgB > GlobalTileTraits
    The traits class to build the iterator to load data from global memory for B^T.
    Definition: igemm_traits.h:289
    gemm_traits.h
    Defines structural properties of complete GEMM computation.
    -
    cutlass::gemm::IgemmTraitsHelper::GlobalLoadStreamB
    GlobalLoadStream< GlobalLoadIteratorB, SharedStoreIteratorB, GlobalTransformerB > GlobalLoadStreamB
    The stream to load B from global memory to shared memory.
    Definition: igemm_traits.h:303
    -
    cutlass::gemm::IgemmTileTraitsHelperA
    Definition: igemm_traits.h:133
    -
    cutlass::gemm::IgemmTraitsHelper::SharedStoreIteratorB
    TileStoreIterator< typename GemmTileTraitsHelperB::SharedStoreTileTraits, typename GemmTileTraitsHelperB::SharedStoreTileTraits::Scalar, IteratorAdvance::kH, MemorySpace::kShared > SharedStoreIteratorB
    The iterator to store B to shared memory.
    Definition: igemm_traits.h:300
    -
    cutlass::gemm::IgemmTraitsHelper::GlobalTransformerB
    IgemmTransformerB< GemmTileTraitsHelperB::kLayout, GlobalLoadIteratorB >::Transformer GlobalTransformerB
    Definition: igemm_traits.h:294
    +
    cutlass::gemm::IgemmTileTraitsHelperA< MatrixLayout::kRowMajor, GemmConfig_, Index_ >::GlobalLoadIterator
    IgemmGlobalIteratorAb< GlobalTileTraits, Index_ > GlobalLoadIterator
    The global load iterator.
    Definition: igemm_traits.h:219
    +
    cutlass::gemm::IgemmTileTraitsHelperA
    Definition: igemm_traits.h:144
    cutlass::gemm::IgemmEpilogue
    Definition: igemm_epilogue.h:290
    -
    cutlass::gemm::IgemmTileTraitsHelperB< MatrixLayout::kRowMajor, GemmConfig_ >::GlobalTileTraits
    IgemmContiguousGlobalTileTraits< GemmOperand::kB, MatrixLayout::kRowMajor, int8_t const, Shape< 1, GemmConfig_::OutputTile::kD, GemmConfig_::OutputTile::kH >, Shape< 1, ShapeCount< typename GemmConfig_::Warps >::kCount, GemmConfig_::kWarpSize >, 4 > GlobalTileTraits
    The traits class to build the iterator to load data from global memory for B^T.
    Definition: igemm_traits.h:203
    cutlass::Copy
    Definition: convert.h:69
    -
    cutlass::gemm::IgemmTileTraitsHelperA< MatrixLayout::kColumnMajor, GemmConfig_ >::Base
    GemmTileTraitsHelperA< MatrixLayout::kColumnMajor, GemmConfig_ > Base
    The base config.
    Definition: igemm_traits.h:141
    -
    cutlass::gemm::IgemmTraitsHelper::GemmConfig
    IgemmConfig< OutputTile_, ScalarD_, AccumulatorsPerThread_ > GemmConfig
    The IGEMM config.
    Definition: igemm_traits.h:267
    +
    cutlass::gemm::IgemmTileTraitsHelperB< MatrixLayout::kRowMajor, GemmConfig_, Index_ >::GlobalTileTraits
    IgemmGlobalTileTraits< GemmOperand::kB, MatrixLayout::kRowMajor, int8_t const, Shape< 1, GemmConfig_::OutputTile::kD, GemmConfig_::OutputTile::kH >, Shape< 1, ShapeCount< typename GemmConfig_::Warps >::kCount, GemmConfig_::kWarpSize >, GemmConfig_::kScalarsPerLdgB > GlobalTileTraits
    The traits class to build the iterator to load data from global memory for B^T.
    Definition: igemm_traits.h:353
    cutlass::gemm::GemmSharedStoreTileAbTraits
    Definition: gemm_shared_tile.h:38
    -
    cutlass::IteratorAdvance::kH
    Definition: tile_iterator.h:62
    +
    cutlass::IteratorAdvance::kH
    Definition: tile_iterator.h:65
    +
    cutlass::gemm::IgemmTileTraitsHelperA< MatrixLayout::kRowMajor, GemmConfig_, Index_ >::MultiplyAddScalar
    int8_t MultiplyAddScalar
    The scalar stored in shared memory.
    Definition: igemm_traits.h:198
    +
    cutlass::gemm::IgemmTraitsHelper::GlobalLoadIteratorB
    GemmTileTraitsHelperB::GlobalLoadIterator GlobalLoadIteratorB
    The iterator to load B from global memory.
    Definition: igemm_traits.h:445
    igemm_multiply_add.h
    Implements matrix multiply accumulate operation of 8-bit integer data using DP4A instruction.
    -
    cutlass::gemm::GemmGlobalIteratorAb
    Definition: gemm_global_tile.h:159
    -
    cutlass::gemm::IgemmTileTraitsHelperB< MatrixLayout::kRowMajor, GemmConfig_ >::SharedStoreTileTraits
    GemmSharedStoreTileAbTraits< int8_t, Shape< GemmConfig_::kStages, GemmConfig_::OutputTile::kD/4, GemmConfig_::OutputTile::kH *4 >, typename GlobalTileTraits::Threads, kScalarsPerStsB > SharedStoreTileTraits
    The traits class to build the iterator to store data to shared memory for B^N.
    Definition: igemm_traits.h:215
    +
    cutlass::gemm::GemmSharedLoadTileBTraits
    Definition: gemm_shared_tile.h:200
    +
    cutlass::gemm::IgemmTraitsHelper::SharedStoreIteratorB
    TileStoreIterator< typename GemmTileTraitsHelperB::SharedStoreTileTraits, typename GemmTileTraitsHelperB::SharedStoreTileTraits::Scalar, IteratorAdvance::kH, MemorySpace::kShared > SharedStoreIteratorB
    The iterator to store B to shared memory.
    Definition: igemm_traits.h:454
    +
    cutlass::gemm::IgemmTileTraitsHelperB< MatrixLayout::kColumnMajor, GemmConfig_, Index_ >::SharedLoadTileTraits
    GemmSharedLoadTileBTraits< int8_t const, typename GemmConfig_::OutputTile, typename GemmConfig_::Warps, typename GemmConfig_::MultiplyAdd::ThreadsPerWarp, typename GemmConfig_::InstructionShape, GemmConfig_::kStages, 16, SharedStoreTileTraits::kSkew > SharedLoadTileTraits
    The traits class to build the iterator to load from shared memory for B^N.
    Definition: igemm_traits.h:326
    +
    cutlass::gemm::GemmGlobalIteratorAb
    Definition: gemm_global_tile.h:163
    +
    cutlass::gemm::IgemmTileTraitsHelperB< MatrixLayout::kColumnMajor, GemmConfig_, Index_ >::MultiplyAddScalar
    int8_t MultiplyAddScalar
    The scalar stored in shared memory.
    Definition: igemm_traits.h:271
    gemm_epilogue.h
    Implements the epilogue phase of the GEMM kernel that efficiently updates global memory with the comp...
    -
    cutlass::gemm::GlobalLoadStream
    Definition: gemm_global_stream.h:161
    -
    cutlass::gemm::GemmTileTraitsHelperB
    Definition: gemm_traits.h:273
    -
    cutlass::gemm::IgemmTraitsHelper::GlobalLoadIteratorB
    GemmGlobalIteratorAb< typename GemmTileTraitsHelperB::GlobalTileTraits, Index_ > GlobalLoadIteratorB
    The iterator to load B from global memory.
    Definition: igemm_traits.h:291
    -
    cutlass::gemm::IgemmTileTraitsHelperA< MatrixLayout::kColumnMajor, GemmConfig_ >::GlobalTileTraits
    IgemmContiguousGlobalTileTraits< GemmOperand::kA, MatrixLayout::kColumnMajor, int8_t const, Shape< 1, GemmConfig_::OutputTile::kD, GemmConfig_::OutputTile::kW >, Shape< 1, ShapeCount< typename GemmConfig_::Warps >::kCount, GemmConfig_::kWarpSize >, 4 > GlobalTileTraits
    The traits class to build the iterator to load data from global memory for A^N.
    Definition: igemm_traits.h:159
    -
    cutlass::gemm::IgemmEpilogueScalar< int >::Scalar
    int Scalar
    Definition: igemm_traits.h:342
    -
    cutlass::gemm::IgemmTransformerB< MatrixLayout::kRowMajor, Iterator_ >::Transformer
    IgemmSwizzle< Iterator_ > Transformer
    Definition: igemm_traits.h:245
    -
    cutlass::MatrixLayout
    Describes layouts of matrices.
    Definition: matrix_traits.h:35
    -
    cutlass::gemm::IgemmTraitsHelper::GemmTileTraitsHelperB
    IgemmTileTraitsHelperB< kLayoutB_, GemmConfig > GemmTileTraitsHelperB
    The GEMM config for B.
    Definition: igemm_traits.h:271
    +
    cutlass::gemm::IgemmTileTraitsHelperA< MatrixLayout::kRowMajor, GemmConfig_, Index_ >::GlobalTileTraits
    IgemmGlobalTileTraits< GemmOperand::kA, MatrixLayout::kRowMajor, int8_t const, Shape< 1, GemmConfig_::OutputTile::kW, GemmConfig_::OutputTile::kD >, Shape< 1, ShapeCount< typename GemmConfig_::Warps >::kCount, GemmConfig_::kWarpSize >, GemmConfig_::kScalarsPerLdgA > GlobalTileTraits
    The traits class to build the iterator to load data from global memory for A^T.
    Definition: igemm_traits.h:216
    +
    cutlass::gemm::GlobalLoadStream
    Definition: gemm_global_stream.h:52
    +
    cutlass::gemm::GemmTileTraitsHelperB
    Definition: gemm_traits.h:191
    +
    cutlass::gemm::IgemmTraitsHelper::Epilogue
    IgemmEpilogue< IgemmEpilogueTraits< GemmConfig, EpilogueFunctor_ > > Epilogue
    The epilogue.
    Definition: igemm_traits.h:487
    +
    cutlass::gemm::IgemmEpilogueScalar< int >::Scalar
    int Scalar
    Definition: igemm_traits.h:499
    +
    cutlass::gemm::IgemmTransformerB< MatrixLayout::kRowMajor, Iterator_ >::Transformer
    IgemmSwizzle< Iterator_ > Transformer
    Definition: igemm_traits.h:398
    +
    cutlass::MatrixLayout
    Defines data layouts of various matrix formats usable by TensorRef and other classes.
    Definition: matrix_traits.h:156
    +
    cutlass::gemm::IgemmTileTraitsHelperA< MatrixLayout::kColumnMajor, GemmConfig_, Index_ >::SharedStoreTileTraits
    GemmSharedStoreTileAbTraits< int8_t, Shape< GemmConfig_::kStages, GemmConfig_::OutputTile::kD/4, GemmConfig_::OutputTile::kW *4 >, typename GlobalTileTraits::Threads, kScalarsPerStsA > SharedStoreTileTraits
    The traits class to build the iterator to store data to shared memory for A^N.
    Definition: igemm_traits.h:185
    cutlass::gemm::IgemmSwizzle
    Definition: igemm_swizzle.h:38
    -
    cutlass::gemm::IgemmTileTraitsHelperB
    Definition: igemm_traits.h:177
    -
    cutlass::gemm::IgemmTraitsHelper
    Definition: igemm_traits.h:265
    -
    cutlass::TileLoadIterator
    An iterator implementing Tile Load Iterator Concept for loading a tile from memory.
    Definition: tile_iterator.h:302
    -
    cutlass::gemm::IgemmTraitsHelper::GlobalLoadStreamA
    GlobalLoadStream< GlobalLoadIteratorA, SharedStoreIteratorA, GlobalTransformerA > GlobalLoadStreamA
    The stream to load A from global memory to shared memory.
    Definition: igemm_traits.h:287
    -
    cutlass::gemm::IgemmTraitsHelper::SharedLoadStreamB
    SharedLoadStream< SharedLoadIteratorB, Copy< typename SharedLoadIteratorB::Fragment > > SharedLoadStreamB
    The stream to load B from shared memory.
    Definition: igemm_traits.h:322
    +
    cutlass::gemm::IgemmTileTraitsHelperB
    Definition: igemm_traits.h:259
    +
    cutlass::gemm::IgemmTraitsHelper
    Definition: igemm_traits.h:418
    +
    cutlass::TileLoadIterator
    An iterator implementing Tile Load Iterator Concept for loading a tile from memory.
    Definition: tile_iterator.h:399
    +
    cutlass::gemm::IgemmTraitsHelper::GlobalTransformerA
    IgemmTransformerA< GemmTileTraitsHelperA::kLayout, GlobalLoadIteratorA >::Transformer GlobalTransformerA
    The default transformer for A.
    Definition: igemm_traits.h:430
    gemm_shared_tile.h
    Defines iterators for efficiently loading and storing tiles to and from shared memory.
    -
    cutlass::MatrixLayout::kRowMajor
    Definition: matrix_traits.h:36
    -
    cutlass::gemm::IgemmTraitsHelper::GemmTileTraitsHelperA
    IgemmTileTraitsHelperA< kLayoutA_, GemmConfig > GemmTileTraitsHelperA
    The GEMM config for A.
    Definition: igemm_traits.h:269
    -
    cutlass::gemm::SharedLoadStream
    Definition: gemm_shared_stream.h:44
    +
    cutlass::gemm::IgemmTraitsHelper::GlobalLoadStreamB
    GlobalLoadStream< GemmOperand::kB, GlobalLoadIteratorB, SharedStoreIteratorB, GlobalTransformerB > GlobalLoadStreamB
    The stream to load B from global memory to shared memory.
    Definition: igemm_traits.h:460
    +
    cutlass::MatrixLayout::kRowMajor
    Definition: matrix_traits.h:159
    +
    cutlass::gemm::SharedLoadStream
    Definition: gemm_shared_stream.h:45
    +
    cutlass::gemm::IgemmGlobalTileTraits
    Definition: igemm_global_tile.h:50
    reshape_tile.h
    Defines a type for restructuring a tile.
    -
    cutlass::gemm::IgemmTraitsHelper::SharedLoadIteratorA
    TileLoadIterator< typename GemmTileTraitsHelperA::SharedLoadTileTraits, typename GemmTileTraitsHelperA::SharedLoadTileTraits::Scalar, IteratorAdvance::kH, MemorySpace::kShared > SharedLoadIteratorA
    The iterator to load A from shared memory.
    Definition: igemm_traits.h:310
    -
    cutlass::gemm::IgemmTraitsHelper::ClearAccumulators
    ClearAccumulators< typename MultiplyAdd::ScalarC > ClearAccumulators
    The object to clear accumulators.
    Definition: igemm_traits.h:327
    -
    cutlass::gemm::GemmConfig
    Definition: gemm_traits.h:79
    -
    cutlass::gemm::GemmTileTraitsHelperA
    Definition: gemm_traits.h:137
    -
    cutlass::GemmOperand::kB
    Definition: matrix_traits.h:43
    +
    cutlass::gemm::IgemmTraitsHelper::GlobalLoadIteratorA
    GemmTileTraitsHelperA::GlobalLoadIterator GlobalLoadIteratorA
    The iterator to load A from global memory.
    Definition: igemm_traits.h:427
    +
    cutlass::gemm::GemmConfig
    Definition: gemm_config.h:76
    +
    cutlass::gemm::IgemmTraitsHelper::SharedStoreIteratorA
    TileStoreIterator< typename GemmTileTraitsHelperA::SharedStoreTileTraits, typename GemmTileTraitsHelperA::SharedStoreTileTraits::Scalar, IteratorAdvance::kH, MemorySpace::kShared > SharedStoreIteratorA
    The iterator to store A to shared memory.
    Definition: igemm_traits.h:436
    +
    cutlass::gemm::GemmTileTraitsHelperA
    Definition: gemm_traits.h:52
    +
    cutlass::GemmOperand::kB
    Definition: matrix_traits.h:357
    cutlass::gemm::IgemmConfig
    Definition: igemm_traits.h:57
    -
    cutlass::gemm::IgemmTransformerA
    Definition: igemm_traits.h:221
    -
    cutlass::gemm::IgemmContiguousGlobalTileTraits
    Definition: igemm_global_tile.h:50
    -
    cutlass::gemm::IgemmEpilogueScalar::Scalar
    float Scalar
    Definition: igemm_traits.h:337
    -
    cutlass::gemm::GemmTraits
    Definition: gemm_traits.h:428
    -
    cutlass::gemm::IgemmTransformerA< MatrixLayout::kRowMajor, Iterator_ >::Transformer
    Copy< typename Iterator_::Fragment > Transformer
    Definition: igemm_traits.h:225
    -
    cutlass::gemm::IgemmTraits
    Definition: igemm_traits.h:370
    +
    cutlass::gemm::IgemmTileTraitsHelperA< MatrixLayout::kColumnMajor, GemmConfig_, Index_ >::GlobalTileTraits
    IgemmGlobalTileTraits< GemmOperand::kA, MatrixLayout::kColumnMajor, int8_t const, Shape< 1, GemmConfig_::OutputTile::kD, GemmConfig_::OutputTile::kW >, Shape< 1, ShapeCount< typename GemmConfig_::Warps >::kCount, GemmConfig_::kWarpSize >, GemmConfig_::kScalarsPerLdgA > GlobalTileTraits
    The traits class to build the iterator to load data from global memory for A^N.
    Definition: igemm_traits.h:170
    +
    cutlass::gemm::IgemmGlobalIteratorAb
    Definition: igemm_global_tile.h:95
    +
    cutlass::gemm::IgemmTransformerA
    Definition: igemm_traits.h:374
    +
    cutlass::gemm::IgemmEpilogueScalar::Scalar
    float Scalar
    Definition: igemm_traits.h:494
    +
    cutlass::gemm::GemmTraits
    Definition: gemm_traits.h:349
    +
    cutlass::gemm::IgemmTransformerA< MatrixLayout::kRowMajor, Iterator_ >::Transformer
    Copy< typename Iterator_::Fragment > Transformer
    Definition: igemm_traits.h:378
    +
    cutlass::gemm::IgemmTraits
    Definition: igemm_traits.h:527
    cutlass::Shape
    A Shape implementing Layout Concept describing the dimensions of a cube.
    Definition: shape.h:64
    -
    cutlass::gemm::IgemmTileTraitsHelperA< MatrixLayout::kColumnMajor, GemmConfig_ >::SharedStoreTileTraits
    GemmSharedStoreTileAbTraits< int8_t, Shape< GemmConfig_::kStages, GemmConfig_::OutputTile::kD/4, GemmConfig_::OutputTile::kW *4 >, typename GlobalTileTraits::Threads, kScalarsPerStsA > SharedStoreTileTraits
    The traits class to build the iterator to store data to shared memory for A^N.
    Definition: igemm_traits.h:171
    -
    cutlass::gemm::GemmTileTraitsHelperB< MatrixLayout::kRowMajor, GemmConfig_ >
    Definition: gemm_traits.h:346
    -
    cutlass::gemm::ThreadMultiplyAdd
    Template performing matrix multiply-add operation within a thread.
    Definition: thread_multiply_add.h:43
    -
    cutlass::MatrixLayout::kColumnMajor
    Definition: matrix_traits.h:36
    -
    cutlass::gemm::GemmTileTraitsHelperA< MatrixLayout::kColumnMajor, GemmConfig_ >
    Definition: gemm_traits.h:142
    -
    cutlass::gemm::IgemmTraitsHelper::Epilogue
    IgemmEpilogue< IgemmEpilogueTraits< GemmConfig, EpilogueFunctor_ > > Epilogue
    The epilogue.
    Definition: igemm_traits.h:330
    -
    cutlass::gemm::IgemmTraitsHelper::GlobalTransformerA
    IgemmTransformerA< GemmTileTraitsHelperA::kLayout, GlobalLoadIteratorA >::Transformer GlobalTransformerA
    The default transformer for A.
    Definition: igemm_traits.h:278
    -
    cutlass::MatrixLayout::Kind
    Kind
    Definition: matrix_traits.h:36
    -
    cutlass::gemm::IgemmTransformerB
    Definition: igemm_traits.h:236
    -
    cutlass::gemm::IgemmTraitsHelper::SharedStoreIteratorA
    TileStoreIterator< typename GemmTileTraitsHelperA::SharedStoreTileTraits, typename GemmTileTraitsHelperA::SharedStoreTileTraits::Scalar, IteratorAdvance::kH, MemorySpace::kShared > SharedStoreIteratorA
    The iterator to store A to shared memory.
    Definition: igemm_traits.h:284
    -
    cutlass::gemm::LinearScaling
    Functor to compute linear combination of fragments.
    Definition: linear_scaling.h:40
    -
    cutlass::GemmOperand::kA
    Definition: matrix_traits.h:43
    +
    cutlass::gemm::IgemmTraitsHelper::SharedLoadIteratorB
    TileLoadIterator< typename GemmTileTraitsHelperB::SharedLoadTileTraits, typename GemmTileTraitsHelperB::SharedLoadTileTraits::Scalar, IteratorAdvance::kH, MemorySpace::kShared > SharedLoadIteratorB
    The iterator to load B from shared memory.
    Definition: igemm_traits.h:476
    +
    cutlass::gemm::IgemmTileTraitsHelperB< MatrixLayout::kRowMajor, GemmConfig_, Index_ >::SharedStoreTileTraits
    GemmSharedStoreTileAbTraits< int8_t, Shape< GemmConfig_::kStages, GemmConfig_::OutputTile::kD/4, GemmConfig_::OutputTile::kH *4 >, typename GlobalTileTraits::Threads, kScalarsPerStsB > SharedStoreTileTraits
    The traits class to build the iterator to store data to shared memory for B^N.
    Definition: igemm_traits.h:368
    +
    cutlass::gemm::GemmGlobalTileTraits::Threads
    ReshapeThreads< VectorizedTile, Threads_ >::Threads Threads
    The threads shape.
    Definition: gemm_global_tile.h:88
    +
    cutlass::gemm::GemmTileTraitsHelperB< MatrixLayout::kRowMajor, GemmConfig_ >
    Definition: gemm_traits.h:267
    +
    cutlass::gemm::ThreadMultiplyAdd
    Template performing matrix multiply-add operation within a thread.
    Definition: thread_multiply_add.h:44
    +
    cutlass::MatrixLayout::kColumnMajor
    Definition: matrix_traits.h:159
    +
    cutlass::gemm::GemmTileTraitsHelperA< MatrixLayout::kColumnMajor, GemmConfig_ >
    Definition: gemm_traits.h:57
    +
    cutlass::gemm::IgemmTileTraitsHelperA< MatrixLayout::kRowMajor, GemmConfig_, Index_ >::Scalar
    int8_t Scalar
    The input scalar.
    Definition: igemm_traits.h:196
    +
    cutlass::gemm::IgemmTraitsHelper::GemmConfig
    IgemmConfig< OutputTile_, ScalarD_, ThreadGemmShape_ > GemmConfig
    The IGEMM config.
    Definition: igemm_traits.h:420
    +
    cutlass::gemm::IgemmTileTraitsHelperB< MatrixLayout::kColumnMajor, GemmConfig_, Index_ >::GlobalLoadIterator
    IgemmGlobalIteratorAb< GlobalTileTraits, Index_ > GlobalLoadIterator
    The global load iterator.
    Definition: igemm_traits.h:292
    +
    cutlass::MatrixLayout::Kind
    Kind
    Enumeration defining fundamental contiguous layouts.
    Definition: matrix_traits.h:159
    +
    cutlass::gemm::IgemmTileTraitsHelperA< MatrixLayout::kColumnMajor, GemmConfig_, Index_ >::GlobalLoadIterator
    GemmGlobalIteratorAb< GlobalTileTraits, Index_ > GlobalLoadIterator
    The global load iterator.
    Definition: igemm_traits.h:173
    +
    cutlass::gemm::IgemmTileTraitsHelperB< MatrixLayout::kRowMajor, GemmConfig_, Index_ >::GlobalLoadIterator
    GemmGlobalIteratorAb< GlobalTileTraits, Index_ > GlobalLoadIterator
    The global load iterator.
    Definition: igemm_traits.h:356
    +
    cutlass::gemm::IgemmTraitsHelper::MultiplyAdd
    GemmConfig::MultiplyAdd MultiplyAdd
    The multiply-add functor.
    Definition: igemm_traits.h:482
    +
    cutlass::gemm::IgemmTransformerB
    Definition: igemm_traits.h:389
    +
    cutlass::gemm::LinearScaling
    Functor to compute linear combination of fragments.
    Definition: linear_scaling.h:51
    +
    cutlass::gemm::IgemmTraitsHelper::SharedLoadStreamA
    SharedLoadStream< SharedLoadIteratorA, Copy< typename SharedLoadIteratorA::Fragment > > SharedLoadStreamA
    The stream to load A from shared memory.
    Definition: igemm_traits.h:470
    +
    cutlass::GemmOperand::kA
    Definition: matrix_traits.h:357
    +
    cutlass::gemm::IgemmTraitsHelper::GlobalLoadStreamA
    GlobalLoadStream< GemmOperand::kA, GlobalLoadIteratorA, SharedStoreIteratorA, GlobalTransformerA > GlobalLoadStreamA
    The stream to load A from global memory to shared memory.
    Definition: igemm_traits.h:442
    +
    cutlass::gemm::IgemmTraitsHelper::GemmTileTraitsHelperA
    IgemmTileTraitsHelperA< kLayoutA_, GemmConfig, Index_ > GemmTileTraitsHelperA
    The GEMM config for A.
    Definition: igemm_traits.h:422
    gemm.h
    Implements a software-pipelined efficient GEMM.
    -
    cutlass::gemm::GemmGlobalTileTraits::Threads
    ReshapeThreads< Tile, Threads_ >::Threads Threads
    The threads shape.
    Definition: gemm_global_tile.h:87
    +
    cutlass::gemm::IgemmTileTraitsHelperA< MatrixLayout::kRowMajor, GemmConfig_, Index_ >::SharedLoadTileTraits
    GemmSharedLoadTileATraits< int8_t const, typename GemmConfig_::OutputTile, typename GemmConfig_::Warps, typename GemmConfig_::MultiplyAdd::ThreadsPerWarp, typename GemmConfig_::InstructionShape, GemmConfig_::kStages, 16, SharedStoreTileTraits::kSkew > SharedLoadTileTraits
    The traits class to build the iterator to load from shared memory for A^N.
    Definition: igemm_traits.h:253
    +
    cutlass::gemm::IgemmTraitsHelper::SharedLoadStreamB
    SharedLoadStream< SharedLoadIteratorB, Copy< typename SharedLoadIteratorB::Fragment > > SharedLoadStreamB
    The stream to load B from shared memory.
    Definition: igemm_traits.h:479
    gemm_epilogue_traits.h
    Defines structural properties of the GEMM epilogue.
    -
    cutlass::gemm::IgemmEpilogueScalar
    Definition: igemm_traits.h:336
    +
    cutlass::gemm::IgemmEpilogueScalar
    Definition: igemm_traits.h:493
    igemm_epilogue.h
    Defines the epilogue phase of the GEMM computation for IGEMM, supporting integer and floating-point o...
    convert.h
    Defines conversion operations among Fragments of different base type.
    -
    cutlass::gemm::IgemmTileTraitsHelperB< MatrixLayout::kRowMajor, GemmConfig_ >::Base
    GemmTileTraitsHelperB< MatrixLayout::kRowMajor, GemmConfig_ > Base
    The base config.
    Definition: igemm_traits.h:185
    -
    cutlass::gemm::IgemmTraitsHelper::SharedLoadStreamA
    SharedLoadStream< SharedLoadIteratorA, Copy< typename SharedLoadIteratorA::Fragment > > SharedLoadStreamA
    The stream to load A from shared memory.
    Definition: igemm_traits.h:313
    +
    cutlass::gemm::IgemmTileTraitsHelperB< MatrixLayout::kColumnMajor, GemmConfig_, Index_ >::SharedStoreTileTraits
    GemmSharedStoreWithSkewTileAbTraits< int8_t, Shape< GemmConfig_::kStages, GemmConfig_::OutputTile::kD/4, GemmConfig_::OutputTile::kH *4 >, typename GlobalTileTraits::Threads, kScalarsPerStsB, 16 > SharedStoreTileTraits
    The traits class to build the iterator to store data to shared memory for B^N.
    Definition: igemm_traits.h:306
    +
    cutlass::gemm::IgemmTileTraitsHelperB< MatrixLayout::kColumnMajor, GemmConfig_, Index_ >::Scalar
    int8_t Scalar
    The input scalar.
    Definition: igemm_traits.h:269
    igemm_global_tile.h
    Implements tile iterators to partition the thread block tile into 2D subtiles and efficiently load ea...
    -
    cutlass::TileStoreIterator
    An iterator implementing Tile Store Iterator Concept for storing a tile to memory.
    Definition: tile_iterator.h:620
    -
    cutlass::gemm::IgemmTraitsHelper::MultiplyAdd
    GemmConfig::MultiplyAdd MultiplyAdd
    The multiply-add functor.
    Definition: igemm_traits.h:325
    +
    cutlass::gemm::IgemmTraitsHelper::SharedLoadIteratorA
    TileLoadIterator< typename GemmTileTraitsHelperA::SharedLoadTileTraits, typename GemmTileTraitsHelperA::SharedLoadTileTraits::Scalar, IteratorAdvance::kH, MemorySpace::kShared > SharedLoadIteratorA
    The iterator to load A from shared memory.
    Definition: igemm_traits.h:467
    +
    cutlass::gemm::IgemmTileTraitsHelperA< MatrixLayout::kColumnMajor, GemmConfig_, Index_ >::Base
    GemmTileTraitsHelperA< MatrixLayout::kColumnMajor, GemmConfig_ > Base
    The base config.
    Definition: igemm_traits.h:152
    +
    cutlass::TileStoreIterator
    An iterator implementing Tile Store Iterator Concept for storing a tile to memory.
    Definition: tile_iterator.h:836
    +
    cutlass::gemm::IgemmTraitsHelper::ClearAccumulators
    ClearAccumulators< typename MultiplyAdd::ScalarC > ClearAccumulators
    The object to clear accumulators.
    Definition: igemm_traits.h:484
    -Generated on Thu May 3 2018 16:36:12 for Cutlass by   +Generated on Wed Sep 19 2018 13:58:50 for Cutlass by   doxygen 1.8.14
    diff --git a/docs/index.html b/docs/index.html index f2ba6899..6fab15e1 100644 --- a/docs/index.html +++ b/docs/index.html @@ -75,7 +75,7 @@ $(function() {
    -Generated on Thu May 3 2018 16:36:15 for Cutlass by   +Generated on Wed Sep 19 2018 13:58:53 for Cutlass by   doxygen 1.8.14
    diff --git a/docs/iterator__access_8h.html b/docs/iterator__access_8h.html index cc41cd5a..06fd90ad 100644 --- a/docs/iterator__access_8h.html +++ b/docs/iterator__access_8h.html @@ -82,10 +82,9 @@ $(function() {

    Free functions for loading and storing to implementations of tile iteartor concepts. More...

    -
    #include <cutlass/fragment_load_store.h>
    -#include <cutlass/load_store.h>
    -#include <cutlass/predicate_vector.h>
    -#include <cutlass/shape.h>
    +
    #include "cutlass/load_store.h"
    +#include "cutlass/predicate_vector.h"
    +#include "cutlass/shape.h"

    Go to the source code of this file.

    @@ -142,7 +146,7 @@ Namespaces

    @@ -98,76 +97,15 @@ Namespaces Functions - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - -
    template<typename InputIterator , typename Fragment >
    CUTLASS_HOST_DEVICE void cutlass::iterator_load (InputIterator &iterator, Fragment &fragment)
     Loads a fragment from an input iterator. More...
     
    template<typename InputIterator , typename Fragment >
    CUTLASS_DEVICE void cutlass::shared_iterator_load (InputIterator &iterator, Fragment &fragment)
     Loads a fragment from a shared memory input iterator. More...
     
    template<typename InputIterator , typename Fragment >
    CUTLASS_DEVICE void cutlass::shared_iterator_load (InputIterator &iterator, Fragment &fragment, int d)
     Loads a fragment from a shared memory input iterator. More...
     
    template<typename InputIterator , typename Fragment , typename ConstPredicateAdapter >
    CUTLASS_HOST_DEVICE void cutlass::iterator_load_post_increment (InputIterator &iterator, Fragment &fragment, typename InputIterator::Index offset, ConstPredicateAdapter predicate_adapter)
     Loads a fragment from an input iterator, masked by a predicate iterator. More...
     
    template<typename InputIterator , typename Fragment >
    CUTLASS_HOST_DEVICE void cutlass::iterator_load_post_increment (InputIterator &iterator, Fragment &fragment, typename InputIterator::Index offset=0)
     Loads a fragment from an input iterator. More...
     
    template<typename InputIterator , typename Fragment , typename ConstPredicateAdapter >
    CUTLASS_HOST_DEVICE void cutlass::iterator_load_post_increment (InputIterator &iterator, Fragment &fragment, ConstPredicateAdapter pred_it)
     Loads a fragment from an input iterator. More...
     
    template<typename InputIterator , typename Fragment , typename ConstPredicateAdapter >
    CUTLASS_HOST_DEVICE void cutlass::iterator_load (InputIterator const &_iterator, Fragment &fragment, typename InputIterator::Index offset, ConstPredicateAdapter predicate_adapter)
     
    template<typename InputIterator , typename Fragment >
    CUTLASS_HOST_DEVICE void cutlass::iterator_load (InputIterator const &iterator, Fragment &fragment, typename InputIterator::Index offset=0)
     Loads a fragment from an input iterator. More...
     
    template<typename InputIterator , typename Fragment , typename ConstPredicateAdapter >
    CUTLASS_HOST_DEVICE void cutlass::iterator_load (InputIterator const &iterator, Fragment &fragment, ConstPredicateAdapter pred_it)
     Loads a fragment from an input iterator. More...
     
    template<typename OutputIterator , typename Fragment >
    CUTLASS_HOST_DEVICE void cutlass::iterator_store (OutputIterator &iterator, Fragment &fragment)
     Stores a fragment to an output iterator. More...
     
    template<typename OutputIterator , typename Fragment >
    CUTLASS_DEVICE void cutlass::shared_iterator_store (OutputIterator &iterator, Fragment const &fragment)
     Stores a fragment to a shared memory output iterator. More...
     
    template<typename OutputIterator , typename Fragment , typename ConstPredicateAdapter >
    CUTLASS_HOST_DEVICE void cutlass::iterator_store_post_increment (OutputIterator &iterator, Fragment const &fragment, typename OutputIterator::Index offset, ConstPredicateAdapter predicate_adapter)
     Stores a fragment to an output iterator, masked by a predicate iterator. More...
     
    template<typename OutputIterator , typename Fragment >
    CUTLASS_HOST_DEVICE void cutlass::iterator_store_post_increment (OutputIterator &iterator, Fragment const &fragment, typename OutputIterator::Index offset=0)
     Stores a fragment to an output iterator. More...
     
    template<typename OutputIterator , typename Fragment , typename ConstPredicateAdapter >
    CUTLASS_HOST_DEVICE void cutlass::iterator_store_post_increment (OutputIterator &iterator, Fragment const &fragment, ConstPredicateAdapter pred_it)
     Stores a fragment to an output iterator. More...
     
    template<typename OutputIterator , typename Fragment , typename ConstPredicateAdapter >
    CUTLASS_HOST_DEVICE void cutlass::iterator_store (OutputIterator const &_iterator, Fragment const &fragment, typename OutputIterator::Index offset, ConstPredicateAdapter predicate_adapter)
     Stores a fragment to an output iterator, masked by a predicate iterator. More...
     
    template<typename OutputIterator , typename Fragment >
    CUTLASS_HOST_DEVICE void cutlass::iterator_store (OutputIterator const &iterator, Fragment const &fragment, typename OutputIterator::Index offset=0)
     Stores a fragment to an output iterator. More...
     
    template<typename OutputIterator , typename Fragment , typename ConstPredicateAdapter >
    CUTLASS_HOST_DEVICE void cutlass::iterator_store (OutputIterator const &iterator, Fragment const &fragment, ConstPredicateAdapter pred_it)
     Stores a fragment to an output iterator. More...
     
    -Generated on Thu May 3 2018 16:36:12 for Cutlass by   +Generated on Wed Sep 19 2018 13:58:50 for Cutlass by   doxygen 1.8.14
    diff --git a/docs/iterator__access_8h_source.html b/docs/iterator__access_8h_source.html index 11289a93..fac9ea1e 100644 --- a/docs/iterator__access_8h_source.html +++ b/docs/iterator__access_8h_source.html @@ -76,30 +76,18 @@ $(function() {
    iterator_access.h
    -Go to the documentation of this file.
    1 /***************************************************************************************************
    2  * Copyright (c) 2017-2018, NVIDIA CORPORATION. All rights reserved.
    3  *
    4  * Redistribution and use in source and binary forms, with or without modification, are permitted
    5  * provided that the following conditions are met:
    6  * * Redistributions of source code must retain the above copyright notice, this list of
    7  * conditions and the following disclaimer.
    8  * * Redistributions in binary form must reproduce the above copyright notice, this list of
    9  * conditions and the following disclaimer in the documentation and/or other materials
    10  * provided with the distribution.
    11  * * Neither the name of the NVIDIA CORPORATION nor the names of its contributors may be used
    12  * to endorse or promote products derived from this software without specific prior written
    13  * permission.
    14  *
    15  * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND ANY EXPRESS OR
    16  * IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND
    17  * FITNESS FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE
    18  * FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING,
    19  * BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS;
    20  * OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT,
    21  * STRICT LIABILITY, OR TOR (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
    22  * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
    23  *
    24  **************************************************************************************************/
    28 #pragma once
    29 
    30 #include <cutlass/fragment_load_store.h>
    31 #include <cutlass/load_store.h>
    32 #include <cutlass/predicate_vector.h>
    33 #include <cutlass/shape.h>
    34 
    35 namespace cutlass {
    36 
    38 
    40 template <typename InputIterator, typename Fragment>
    41 CUTLASS_HOST_DEVICE void iterator_load(InputIterator &iterator, Fragment &fragment) {
    42  typename InputIterator::FragmentIterator frag_iterator(fragment);
    43  for (int d = 0; d < InputIterator::Iterations::kD; ++d) {
    44  for (int h = 0; h < InputIterator::Iterations::kH; ++h) {
    45  for (int w = 0; w < InputIterator::Iterations::kW; ++w) {
    46  for (int c = 0; c < InputIterator::Iterations::kC; ++c) {
    47  if (iterator.valid(d, h, w, c)) {
    48  int const offset =
    49  ComputeOffsetFromStrides<typename InputIterator::ImmediateOffsetStrides>::get(
    50  0, 0, w, c);
    51  Load<typename Fragment::Element, InputIterator::Tile::kC, InputIterator::kMemorySpace>::
    52  load(reinterpret_cast<typename InputIterator::AccessType &>(
    53  frag_iterator.at(d, h, w, c)),
    54  iterator.data(),
    55  offset);
    56  }
    57  }
    58  if (w < InputIterator::Iterations::kW - 1) {
    59  iterator.inc_w();
    60  }
    61  }
    62  if (h < InputIterator::Iterations::kH - 1) {
    63  iterator.inc_h();
    64  }
    65  }
    66  if (d < InputIterator::Iterations::kD - 1) {
    67  iterator.inc_d();
    68  }
    69  }
    70  iterator.inc_advance();
    71 }
    72 
    74 template <typename InputIterator, typename Fragment>
    75 CUTLASS_DEVICE void shared_iterator_load(InputIterator &iterator, Fragment &fragment) {
    76  typename InputIterator::FragmentIterator frag_iterator(fragment);
    77  for (int d = 0; d < InputIterator::Iterations::kD; ++d) {
    78  for (int h = 0; h < InputIterator::Iterations::kH; ++h) {
    79  for (int w = 0; w < InputIterator::Iterations::kW; ++w) {
    80  for (int c = 0; c < InputIterator::Iterations::kC; ++c) {
    81  int const offset =
    82  ComputeOffsetFromStrides<typename InputIterator::ImmediateOffsetStrides>::get(
    83  d, h, w, c);
    84 
    85  FragmentLoad<InputIterator::kIteratorFragment,
    86  InputIterator::Tile::kC,
    87  typename InputIterator::Scalar,
    88  InputIterator::kMemorySpace,
    89  typename InputIterator::FragmentElement,
    90  InputIterator::Tile::kW>::load(frag_iterator.at(d, h, w, c),
    91  iterator.data(),
    92  offset);
    93  }
    94  }
    95  }
    96  }
    97 }
    98 
    100 template <typename InputIterator, typename Fragment>
    101 CUTLASS_DEVICE void shared_iterator_load(InputIterator &iterator, Fragment &fragment, int d) {
    102  typename InputIterator::FragmentIterator frag_iterator(fragment);
    103  for (int h = 0; h < InputIterator::Iterations::kH; ++h) {
    104  for (int w = 0; w < InputIterator::Iterations::kW; ++w) {
    105  for (int c = 0; c < InputIterator::Iterations::kC; ++c) {
    106  int const offset =
    107  ComputeOffsetFromStrides<typename InputIterator::ImmediateOffsetStrides>::get(
    108  d, h, w, c);
    109 
    110  FragmentLoad<InputIterator::kIteratorFragment,
    111  InputIterator::Tile::kC,
    112  typename InputIterator::Scalar,
    113  InputIterator::kMemorySpace,
    114  typename InputIterator::FragmentElement,
    115  InputIterator::Tile::kW>::load(frag_iterator.at(0, h, w, c),
    116  iterator.data(),
    117  offset);
    118  }
    119  }
    120  }
    121 }
    122 
    124 template <typename InputIterator, typename Fragment, typename ConstPredicateAdapter>
    125 CUTLASS_HOST_DEVICE void iterator_load_post_increment(InputIterator &iterator,
    126  Fragment &fragment,
    127  typename InputIterator::Index offset,
    128  ConstPredicateAdapter predicate_adapter) {
    129  for (int d = 0; d < InputIterator::Iterations::kD; ++d, iterator.inc_d()) {
    130  for (int h = 0; h < InputIterator::Iterations::kH; ++h, iterator.inc_h()) {
    131  for (int w = 0; w < InputIterator::Iterations::kW; ++w, iterator.inc_w()) {
    132  if (predicate_adapter.at(d, h, w, 0)) {
    133  int idx = InputIterator::Tile::kC *
    134  (w + InputIterator::Iterations::kW * (h + InputIterator::Iterations::kH * d));
    135 
    136  Load<typename Fragment::Element, InputIterator::Tile::kC, InputIterator::kMemorySpace>::
    137  load(reinterpret_cast<typename InputIterator::AccessType &>(fragment[idx]),
    138  iterator.data(),
    139  offset);
    140  }
    141  }
    142  }
    143  }
    144 }
    145 
    147 template <typename InputIterator, typename Fragment>
    148 CUTLASS_HOST_DEVICE void iterator_load_post_increment(InputIterator &iterator,
    149  Fragment &fragment,
    150  typename InputIterator::Index offset = 0) {
    151  TrivialPredicateTileAdapter pred;
    152  iterator_load_post_increment(iterator, fragment, offset, pred);
    153 }
    154 
    156 template <typename InputIterator, typename Fragment, typename ConstPredicateAdapter>
    157 CUTLASS_HOST_DEVICE void iterator_load_post_increment(InputIterator &iterator,
    158  Fragment &fragment,
    159  ConstPredicateAdapter pred_it) {
    160  iterator_load_post_increment(iterator, fragment, 0, pred_it);
    161 }
    162 
    163 template <typename InputIterator, typename Fragment, typename ConstPredicateAdapter>
    164 CUTLASS_HOST_DEVICE void iterator_load(InputIterator const &_iterator,
    165  Fragment &fragment,
    166  typename InputIterator::Index offset,
    167  ConstPredicateAdapter predicate_adapter) {
    168  InputIterator iterator(_iterator);
    169  iterator_load_post_increment(iterator, fragment, offset, predicate_adapter);
    170 }
    171 
    173 template <typename InputIterator, typename Fragment>
    174 CUTLASS_HOST_DEVICE void iterator_load(InputIterator const &iterator,
    175  Fragment &fragment,
    176  typename InputIterator::Index offset = 0) {
    177  TrivialPredicateTileAdapter pred;
    178  iterator_load(iterator, fragment, offset, pred);
    179 }
    180 
    182 template <typename InputIterator, typename Fragment, typename ConstPredicateAdapter>
    183 CUTLASS_HOST_DEVICE void iterator_load(InputIterator const &iterator,
    184  Fragment &fragment,
    185  ConstPredicateAdapter pred_it) {
    186  iterator_load(iterator, fragment, 0, pred_it);
    187 }
    188 
    190 
    192 template <typename OutputIterator, typename Fragment>
    193 CUTLASS_HOST_DEVICE void iterator_store(OutputIterator &iterator, Fragment &fragment) {
    194  typename OutputIterator::FragmentIterator frag_iterator(fragment);
    195  for (int d = 0; d < OutputIterator::Iterations::kD; ++d) {
    196  for (int h = 0; h < OutputIterator::Iterations::kH; ++h) {
    197  for (int w = 0; w < OutputIterator::Iterations::kW; ++w) {
    198  if (iterator.valid(d, h, w, 0)) {
    199  int const offset =
    200  ComputeOffsetFromStrides<typename OutputIterator::ImmediateOffsetStrides>::get(
    201  d, h, w, 0);
    202 
    203  Store<typename Fragment::Element,
    204  OutputIterator::Tile::kC,
    205  OutputIterator::kMemorySpace>::
    206  store(reinterpret_cast<typename OutputIterator::AccessType &>(
    207  frag_iterator.at(d, h, w, 0)),
    208  iterator.data(),
    209  offset);
    210  }
    211  if (w < OutputIterator::Iterations::kW - 1) {
    212  iterator.inc_w();
    213  }
    214  }
    215  if (h < OutputIterator::Iterations::kH - 1) {
    216  iterator.inc_h();
    217  }
    218  }
    219  if (d < OutputIterator::Iterations::kD - 1) {
    220  iterator.inc_d();
    221  }
    222  }
    223  iterator.inc_advance();
    224 }
    225 
    227 template <typename OutputIterator, typename Fragment>
    228 CUTLASS_DEVICE void shared_iterator_store(OutputIterator &iterator, Fragment const &fragment) {
    229  typename OutputIterator::FragmentConstIterator frag_iterator(fragment);
    230  for (int d = 0; d < OutputIterator::Iterations::kD; ++d) {
    231  for (int h = 0; h < OutputIterator::Iterations::kH; ++h) {
    232  for (int w = 0; w < OutputIterator::Iterations::kW; ++w) {
    233  for (int c = 0; c < OutputIterator::Iterations::kC; ++c) {
    234  int const offset =
    235  ComputeOffsetFromStrides<typename OutputIterator::ImmediateOffsetStrides>::get(
    236  d, h, w, c);
    237 
    238  FragmentStore<OutputIterator::kIteratorFragment,
    239  OutputIterator::Tile::kC,
    240  typename OutputIterator::Scalar,
    241  OutputIterator::kMemorySpace,
    242  typename OutputIterator::FragmentElement,
    243  OutputIterator::Tile::kW>::store(frag_iterator.at(d, h, w, c),
    244  iterator.data(),
    245  offset);
    246  }
    247  }
    248  }
    249  }
    250 }
    251 
    253 
    255 template <typename OutputIterator, typename Fragment, typename ConstPredicateAdapter>
    256 CUTLASS_HOST_DEVICE void iterator_store_post_increment(OutputIterator &iterator,
    257  Fragment const &fragment,
    258  typename OutputIterator::Index offset,
    259  ConstPredicateAdapter predicate_adapter) {
    260  for (int d = 0; d < OutputIterator::Iterations::kD; ++d, iterator.inc_d()) {
    261  for (int h = 0; h < OutputIterator::Iterations::kH; ++h, iterator.inc_h()) {
    262  for (int w = 0; w < OutputIterator::Iterations::kW; ++w, iterator.inc_w()) {
    263  if (predicate_adapter.at(d, h, w, 0)) {
    264  int idx = OutputIterator::Tile::kC *
    265  (w + OutputIterator::Iterations::kW * (h + OutputIterator::Iterations::kH * d));
    266 
    267  Store<typename Fragment::Element,
    268  OutputIterator::Tile::kC,
    269  OutputIterator::kMemorySpace>::
    270  store(reinterpret_cast<typename OutputIterator::AccessType const &>(fragment[idx]),
    271  iterator.data(),
    272  offset);
    273  }
    274  }
    275  }
    276  }
    277 }
    278 
    280 template <typename OutputIterator, typename Fragment>
    281 CUTLASS_HOST_DEVICE void iterator_store_post_increment(OutputIterator &iterator,
    282  Fragment const &fragment,
    283  typename OutputIterator::Index offset = 0) {
    284  TrivialPredicateTileAdapter pred;
    285  iterator_store_post_increment(iterator, fragment, offset, pred);
    286 }
    287 
    289 template <typename OutputIterator, typename Fragment, typename ConstPredicateAdapter>
    290 CUTLASS_HOST_DEVICE void iterator_store_post_increment(OutputIterator &iterator,
    291  Fragment const &fragment,
    292  ConstPredicateAdapter pred_it) {
    293  iterator_store_post_increment(iterator, fragment, 0, pred_it);
    294 }
    295 
    297 template <typename OutputIterator, typename Fragment, typename ConstPredicateAdapter>
    298 CUTLASS_HOST_DEVICE void iterator_store(OutputIterator const &_iterator,
    299  Fragment const &fragment,
    300  typename OutputIterator::Index offset,
    301  ConstPredicateAdapter predicate_adapter) {
    302  OutputIterator iterator(_iterator);
    303  iterator_store_post_increment(iterator, fragment, offset, predicate_adapter);
    304 }
    305 
    307 template <typename OutputIterator, typename Fragment>
    308 CUTLASS_HOST_DEVICE void iterator_store(OutputIterator const &iterator,
    309  Fragment const &fragment,
    310  typename OutputIterator::Index offset = 0) {
    311  TrivialPredicateTileAdapter pred;
    312  iterator_store(iterator, fragment, offset, pred);
    313 }
    314 
    316 template <typename OutputIterator, typename Fragment, typename ConstPredicateAdapter>
    317 CUTLASS_HOST_DEVICE void iterator_store(OutputIterator const &iterator,
    318  Fragment const &fragment,
    319  ConstPredicateAdapter pred_it) {
    320  iterator_store(iterator, fragment, 0, pred_it);
    321 }
    322 
    324 
    325 } // namespace cutlass
    cutlass::FragmentLoad
    Definition: fragment_load_store.h:43
    -
    cutlass
    Definition: convert.h:33
    -
    cutlass::shared_iterator_load
    CUTLASS_DEVICE void shared_iterator_load(InputIterator &iterator, Fragment &fragment)
    Loads a fragment from a shared memory input iterator.
    Definition: iterator_access.h:75
    -
    cutlass::iterator_store_post_increment
    CUTLASS_HOST_DEVICE void iterator_store_post_increment(OutputIterator &iterator, Fragment const &fragment, typename OutputIterator::Index offset, ConstPredicateAdapter predicate_adapter)
    Stores a fragment to an output iterator, masked by a predicate iterator.
    Definition: iterator_access.h:256
    -
    fragment_load_store.h
    Defines accessors for loading and storing fragments to memory efficiently.
    -
    cutlass::Load::load
    static CUTLASS_DEVICE void load(AccessType &dst, Scalar_ const *pointer, int offset)
    The load function.
    Definition: load_store.h:59
    +Go to the documentation of this file.
    1 /***************************************************************************************************
    2  * Copyright (c) 2017-2018, NVIDIA CORPORATION. All rights reserved.
    3  *
    4  * Redistribution and use in source and binary forms, with or without modification, are permitted
    5  * provided that the following conditions are met:
    6  * * Redistributions of source code must retain the above copyright notice, this list of
    7  * conditions and the following disclaimer.
    8  * * Redistributions in binary form must reproduce the above copyright notice, this list of
    9  * conditions and the following disclaimer in the documentation and/or other materials
    10  * provided with the distribution.
    11  * * Neither the name of the NVIDIA CORPORATION nor the names of its contributors may be used
    12  * to endorse or promote products derived from this software without specific prior written
    13  * permission.
    14  *
    15  * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND ANY EXPRESS OR
    16  * IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND
    17  * FITNESS FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE
    18  * FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING,
    19  * BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS;
    20  * OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT,
    21  * STRICT LIABILITY, OR TOR (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
    22  * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
    23  *
    24  **************************************************************************************************/
    28 #pragma once
    29 
    30 #include "cutlass/load_store.h"
    31 #include "cutlass/predicate_vector.h"
    32 #include "cutlass/shape.h"
    33 
    34 namespace cutlass {
    35 
    37 // Used by convolution
    38 template <typename InputIterator, typename Fragment>
    39 CUTLASS_HOST_DEVICE void iterator_load(InputIterator &iterator, Fragment &fragment) {
    40  typename InputIterator::FragmentIterator frag_iterator(fragment);
    41  for (int d = 0; d < InputIterator::Iterations::kD; ++d) {
    42  for (int h = 0; h < InputIterator::Iterations::kH; ++h) {
    43  for (int w = 0; w < InputIterator::Iterations::kW; ++w) {
    44  for (int c = 0; c < InputIterator::Iterations::kC; ++c) {
    45  if (iterator.valid(d, h, w, c)) {
    46  iterator.load_element(reinterpret_cast<typename InputIterator::AccessType &>(
    47  frag_iterator.at(d, h, w, c)),
    48  d,
    49  h,
    50  w,
    51  c);
    52  }
    53  }
    54  if (w < InputIterator::Iterations::kW - 1) {
    55  iterator.inc_w();
    56  }
    57  }
    58  if (h < InputIterator::Iterations::kH - 1) {
    59  iterator.inc_h();
    60  }
    61  }
    62  if (d < InputIterator::Iterations::kD - 1) {
    63  iterator.inc_d();
    64  }
    65  }
    66  iterator.inc_advance();
    67 }
    68 
    69 template <typename OutputIterator, typename Fragment>
    70 CUTLASS_HOST_DEVICE void iterator_store(OutputIterator &iterator, Fragment &fragment) {
    71  typename OutputIterator::FragmentIterator frag_iterator(fragment);
    72  for (int d = 0; d < OutputIterator::Iterations::kD; ++d) {
    73  for (int h = 0; h < OutputIterator::Iterations::kH; ++h) {
    74  for (int w = 0; w < OutputIterator::Iterations::kW; ++w) {
    75  for (int c = 0; c < OutputIterator::Iterations::kC; ++c) {
    76  if (iterator.valid(d, h, w, c)) {
    77  iterator.store_element(reinterpret_cast<typename OutputIterator::AccessType &>(
    78  frag_iterator.at(d, h, w, c)),
    79  d,
    80  h,
    81  w,
    82  c);
    83  }
    84  }
    85  if (w < OutputIterator::Iterations::kW - 1) {
    86  iterator.inc_w();
    87  }
    88  }
    89  if (h < OutputIterator::Iterations::kH - 1) {
    90  iterator.inc_h();
    91  }
    92  }
    93  if (d < OutputIterator::Iterations::kD - 1) {
    94  iterator.inc_d();
    95  }
    96  }
    97  iterator.inc_advance();
    98 }
    100 
    101 } // namespace cutlass
    cutlass
    Definition: convert.h:33
    cutlass::Fragment
    A template defining Fragment Concept.
    Definition: fragment.h:99
    -
    cutlass::Store
    Definition: load_store.h:131
    predicate_vector.h
    Defines container classes and iterators for managing a statically sized vector of boolean predicates...
    -
    cutlass::ComputeOffsetFromStrides::get
    static CUTLASS_DEVICE int get(int d, int h, int w, int c)
    Definition: shape.h:211
    -
    cutlass::iterator_load_post_increment
    CUTLASS_HOST_DEVICE void iterator_load_post_increment(InputIterator &iterator, Fragment &fragment, typename InputIterator::Index offset, ConstPredicateAdapter predicate_adapter)
    Loads a fragment from an input iterator, masked by a predicate iterator.
    Definition: iterator_access.h:125
    load_store.h
    Defines abstractions for efficiently loading and storing vectors to memory.
    CUTLASS_HOST_DEVICE
    #define CUTLASS_HOST_DEVICE
    Definition: cutlass.h:46
    -
    cutlass::shared_iterator_store
    CUTLASS_DEVICE void shared_iterator_store(OutputIterator &iterator, Fragment const &fragment)
    Stores a fragment to a shared memory output iterator.
    Definition: iterator_access.h:228
    -
    cutlass::Fragment::Element
    Element_ Element
    The element.
    Definition: fragment.h:108
    -
    cutlass::TrivialPredicateTileAdapter
    Always returns true predicate.
    Definition: predicate_vector.h:426
    -
    cutlass::iterator_store
    CUTLASS_HOST_DEVICE void iterator_store(OutputIterator &iterator, Fragment &fragment)
    Stores a fragment to an output iterator.
    Definition: iterator_access.h:193
    -
    cutlass::FragmentStore
    Definition: fragment_load_store.h:91
    -
    cutlass::iterator_load
    CUTLASS_HOST_DEVICE void iterator_load(InputIterator &iterator, Fragment &fragment)
    Loads a fragment from an input iterator.
    Definition: iterator_access.h:41
    +
    cutlass::iterator_store
    CUTLASS_HOST_DEVICE void iterator_store(OutputIterator &iterator, Fragment &fragment)
    Definition: iterator_access.h:70
    +
    cutlass::iterator_load
    CUTLASS_HOST_DEVICE void iterator_load(InputIterator &iterator, Fragment &fragment)
    Definition: iterator_access.h:39
    shape.h
    Defines Shape implementing the Layout concept for representing a 4D hypercube of objects.
    -Generated on Thu May 3 2018 16:36:12 for Cutlass by   +Generated on Wed Sep 19 2018 13:58:50 for Cutlass by   doxygen 1.8.14
    diff --git a/docs/kernel__launch_8h.html b/docs/kernel__launch_8h.html new file mode 100644 index 00000000..192d541b --- /dev/null +++ b/docs/kernel__launch_8h.html @@ -0,0 +1,108 @@ + + + + + + + +Cutlass: kernel_launch.h File Reference + + + + + + + + + + +
    +
    + + + + + + +
    +
    Cutlass +
    +
    CUDA Templates for Linear Algebra Subroutines and Solvers
    +
    +
    + + + + + + + + +
    +
    + + +
    + +
    + + +
    +
    +
    +Classes | +Namespaces
    +
    +
    kernel_launch.h File Reference
    +
    +
    + +

    Defines structures and helpers to launch CUDA kernels within CUTLASS. +More...

    +
    #include "cutlass/cutlass.h"
    +
    +

    Go to the source code of this file.

    + + + + + +

    +Classes

    struct  cutlass::KernelLaunchConfiguration
     Structure containing the basic launch configuration of a CUDA kernel. More...
     
    + + + +

    +Namespaces

     cutlass
     
    +
    + +
    +Generated on Wed Sep 19 2018 13:58:50 for Cutlass by   +doxygen + 1.8.14 +
    + + diff --git a/docs/kernel__launch_8h_source.html b/docs/kernel__launch_8h_source.html new file mode 100644 index 00000000..52c7a5e0 --- /dev/null +++ b/docs/kernel__launch_8h_source.html @@ -0,0 +1,95 @@ + + + + + + + +Cutlass: kernel_launch.h Source File + + + + + + + + + + +
    +
    + + + + + + +
    +
    Cutlass +
    +
    CUDA Templates for Linear Algebra Subroutines and Solvers
    +
    +
    + + + + + + + + +
    +
    + + +
    + +
    + + +
    +
    +
    +
    kernel_launch.h
    +
    +
    +Go to the documentation of this file.
    1 /***************************************************************************************************
    2  * Copyright (c) 2017-2018, NVIDIA CORPORATION. All rights reserved.
    3  *
    4  * Redistribution and use in source and binary forms, with or without modification, are permitted
    5  * provided that the following conditions are met:
    6  * * Redistributions of source code must retain the above copyright notice, this list of
    7  * conditions and the following disclaimer.
    8  * * Redistributions in binary form must reproduce the above copyright notice, this list of
    9  * conditions and the following disclaimer in the documentation and/or other materials
    10  * provided with the distribution.
    11  * * Neither the name of the NVIDIA CORPORATION nor the names of its contributors may be used
    12  * to endorse or promote products derived from this software without specific prior written
    13  * permission.
    14  *
    15  * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND ANY EXPRESS OR
    16  * IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND
    17  * FITNESS FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE
    18  * FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING,
    19  * BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS;
    20  * OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT,
    21  * STRICT LIABILITY, OR TOR (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
    22  * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
    23  *
    24  **************************************************************************************************/
    29 #pragma once
    30 
    31 #include "cutlass/cutlass.h"
    32 
    33 namespace cutlass {
    34 
    36 
    38 struct KernelLaunchConfiguration {
    39 
    41  dim3 grid;
    42 
    44  dim3 block;
    45 
    47  size_t dynamic_smem;
    48 
    49  //
    50  // Methods
    51  //
    52 
    54  CUTLASS_HOST_DEVICE
    55  KernelLaunchConfiguration(
    56  dim3 _grid = dim3(1,1,1),
    57  dim3 _block = dim3(1,1,1),
    58  size_t _dynamic_smem = 0
    59  ):
    60  grid(_grid),
    61  block(_block),
    62  dynamic_smem(_dynamic_smem) { }
    63 };
    64 
    66 
    67 } // namespace cutlass
    cutlass::KernelLaunchConfiguration::KernelLaunchConfiguration
    CUTLASS_HOST_DEVICE KernelLaunchConfiguration(dim3 _grid=dim3(1, 1, 1), dim3 _block=dim3(1, 1, 1), size_t _dynamic_smem=0)
    Constructs a KernellaunchConfiguration object.
    Definition: kernel_launch.h:55
    +
    cutlass
    Definition: convert.h:33
    +
    cutlass::KernelLaunchConfiguration
    Structure containing the basic launch configuration of a CUDA kernel.
    Definition: kernel_launch.h:38
    +
    CUTLASS_HOST_DEVICE
    #define CUTLASS_HOST_DEVICE
    Definition: cutlass.h:46
    +
    cutlass::KernelLaunchConfiguration::dynamic_smem
    size_t dynamic_smem
    Bytes of dynamically allocated SMEM in addition to static SMEM.
    Definition: kernel_launch.h:47
    +
    cutlass::KernelLaunchConfiguration::block
    dim3 block
    CUDA threablock dimensions.
    Definition: kernel_launch.h:44
    +
    cutlass::KernelLaunchConfiguration::grid
    dim3 grid
    CUDA grid dimensions.
    Definition: kernel_launch.h:41
    +
    cutlass.h
    Basic include for CUTLASS macros.
    +
    + +
    +Generated on Wed Sep 19 2018 13:58:50 for Cutlass by   +doxygen + 1.8.14 +
    + + diff --git a/docs/linear__scaling_8h.html b/docs/linear__scaling_8h.html index 060be3aa..132c09d4 100644 --- a/docs/linear__scaling_8h.html +++ b/docs/linear__scaling_8h.html @@ -74,7 +74,8 @@ $(function() {
    Classes | -Namespaces
    +Namespaces | +Functions
    linear_scaling.h File Reference
    @@ -82,7 +83,7 @@ $(function() {

    Implements the BLAS linear scaling function alpha*AB + beta*C. More...

    -
    #include <cutlass/fragment_multiply_add.h>
    +
    #include "cutlass/fragment_multiply_add.h"

    Go to the source code of this file.

    @@ -101,11 +102,19 @@ Namespaces +
     
     cutlass::gemm
     
    + + + + + +

    +Functions

    template<typename T >
    CUTLASS_DEVICE bool cutlass::gemm::is_zero (T x)
     
    CUTLASS_DEVICE bool cutlass::gemm::is_zero (half x)
     
    -Generated on Thu May 3 2018 16:36:12 for Cutlass by   +Generated on Wed Sep 19 2018 13:58:50 for Cutlass by   doxygen 1.8.14
    diff --git a/docs/linear__scaling_8h_source.html b/docs/linear__scaling_8h_source.html index d9817ed0..b00e5859 100644 --- a/docs/linear__scaling_8h_source.html +++ b/docs/linear__scaling_8h_source.html @@ -76,25 +76,33 @@ $(function() {
    linear_scaling.h
    -Go to the documentation of this file.
    1 
    2 /***************************************************************************************************
    3  * Copyright (c) 2017-2018, NVIDIA CORPORATION. All rights reserved.
    4  *
    5  * Redistribution and use in source and binary forms, with or without modification, are permitted
    6  * provided that the following conditions are met:
    7  * * Redistributions of source code must retain the above copyright notice, this list of
    8  * conditions and the following disclaimer.
    9  * * Redistributions in binary form must reproduce the above copyright notice, this list of
    10  * conditions and the following disclaimer in the documentation and/or other materials
    11  * provided with the distribution.
    12  * * Neither the name of the NVIDIA CORPORATION nor the names of its contributors may be used
    13  * to endorse or promote products derived from this software without specific prior written
    14  * permission.
    15  *
    16  * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND ANY EXPRESS OR
    17  * IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND
    18  * FITNESS FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE
    19  * FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING,
    20  * BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS;
    21  * OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT,
    22  * STRICT LIABILITY, OR TOR (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
    23  * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
    24  *
    25  **************************************************************************************************/
    29 #pragma once
    30 
    31 #include <cutlass/fragment_multiply_add.h>
    32 
    33 namespace cutlass {
    34 namespace gemm {
    35 
    37 
    39 template <typename Scalar_, typename FragmentMultiplyAdd_ = FragmentMultiplyAdd<Scalar_> >
    40 struct LinearScaling {
    41  // The scalar.
    42  typedef Scalar_ Scalar;
    43  // The adapater.
    44  typedef FragmentMultiplyAdd_ FragmentMultiplyAdd;
    45 
    47  struct Params {
    49  Scalar alpha, beta;
    50 
    52  template <typename GemmDesc_>
    53  CUTLASS_HOST_DEVICE int initialize(GemmDesc_ const& desc) {
    54  alpha = desc.alpha;
    55  beta = desc.beta;
    56  return 0;
    57  }
    58  };
    59 
    61  CUTLASS_DEVICE LinearScaling(Params const& params) : alpha(params.alpha), beta(params.beta) {}
    62 
    64  template <typename Fragment_>
    65  CUTLASS_DEVICE void evaluate(Fragment_ const& accum, Fragment_& output) {
    66  FragmentMultiplyAdd mad;
    67  mad.multiply(alpha, accum, output);
    68  }
    69 
    71  template <typename Fragment_>
    72  CUTLASS_DEVICE void evaluate(Fragment_ const& accum, Fragment_ const& old, Fragment_& output) {
    73  FragmentMultiplyAdd mad;
    74  Fragment_ tmp;
    75  mad.multiply(beta, old, tmp);
    76  mad.multiply_add(alpha, accum, tmp, output);
    77  }
    78 
    80  Scalar alpha, beta;
    81 };
    82 
    84 
    85 } // namespace gemm
    86 } // namespace cutlass
    cutlass
    Definition: convert.h:33
    -
    cutlass::gemm::LinearScaling::Params::alpha
    Scalar alpha
    The alpha/beta scaling params.
    Definition: linear_scaling.h:49
    -
    cutlass::gemm::LinearScaling::alpha
    Scalar alpha
    The alpha/beta scaling factors.
    Definition: linear_scaling.h:80
    -
    cutlass::gemm::LinearScaling::LinearScaling
    CUTLASS_DEVICE LinearScaling(Params const &params)
    Ctor.
    Definition: linear_scaling.h:61
    -
    cutlass::gemm::LinearScaling::evaluate
    CUTLASS_DEVICE void evaluate(Fragment_ const &accum, Fragment_ const &old, Fragment_ &output)
    Evaluate the functor.
    Definition: linear_scaling.h:72
    -
    cutlass::gemm::LinearScaling::Params::beta
    Scalar beta
    Definition: linear_scaling.h:49
    -
    cutlass::gemm::LinearScaling::Params::initialize
    CUTLASS_HOST_DEVICE int initialize(GemmDesc_ const &desc)
    Initialize the parameters.
    Definition: linear_scaling.h:53
    -
    cutlass::gemm::LinearScaling::beta
    Scalar beta
    Definition: linear_scaling.h:80
    +Go to the documentation of this file.
    1 
    2 /***************************************************************************************************
    3  * Copyright (c) 2017-2018, NVIDIA CORPORATION. All rights reserved.
    4  *
    5  * Redistribution and use in source and binary forms, with or without modification, are permitted
    6  * provided that the following conditions are met:
    7  * * Redistributions of source code must retain the above copyright notice, this list of
    8  * conditions and the following disclaimer.
    9  * * Redistributions in binary form must reproduce the above copyright notice, this list of
    10  * conditions and the following disclaimer in the documentation and/or other materials
    11  * provided with the distribution.
    12  * * Neither the name of the NVIDIA CORPORATION nor the names of its contributors may be used
    13  * to endorse or promote products derived from this software without specific prior written
    14  * permission.
    15  *
    16  * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND ANY EXPRESS OR
    17  * IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND
    18  * FITNESS FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE
    19  * FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING,
    20  * BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS;
    21  * OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT,
    22  * STRICT LIABILITY, OR TOR (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
    23  * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
    24  *
    25  **************************************************************************************************/
    29 #pragma once
    30 
    31 #include "cutlass/fragment_multiply_add.h"
    32 
    33 namespace cutlass {
    34 namespace gemm {
    35 
    37 
    38 template <typename T>
    39 CUTLASS_DEVICE bool is_zero(T x) {
    40  return x == T(0);
    41 }
    42 
    43 #if !defined(__CUDACC_RTC__) || defined(CUTLASS_NVRTC_HAS_FP16)
    44 CUTLASS_DEVICE bool is_zero(half x) { return reinterpret_cast<int16_t&>(x) == int16_t(0); }
    45 #endif
    46 
    48 
    50 template <typename Scalar_, typename FragmentMultiplyAdd_ = FragmentMultiplyAdd<Scalar_, Scalar_> >
    51 struct LinearScaling {
    52  // The scalar.
    53  typedef Scalar_ Scalar;
    54  // The accumulator Type
    55  typedef typename FragmentMultiplyAdd_::ScalarAccum ScalarAccum;
    56  // The adapater.
    57  typedef FragmentMultiplyAdd_ FragmentMultiplyAdd;
    58 
    60  struct Params {
    62  Scalar alpha, beta;
    63 
    64  //
    65  // Methods
    66  //
    67 
    68  // Constructor
    69  CUTLASS_HOST_DEVICE
    70  Params(Scalar _alpha = 0, Scalar _beta = 0) : alpha(_alpha), beta(_beta) {}
    71 
    73  CUTLASS_HOST_DEVICE int initialize(Scalar _alpha, Scalar _beta) {
    74  alpha = _alpha;
    75  beta = _beta;
    76  return 0;
    77  }
    78 
    80  template <typename GemmDesc_>
    81  CUTLASS_HOST_DEVICE int initialize(GemmDesc_ const& desc) {
    82  alpha = desc.alpha;
    83  beta = desc.beta;
    84  return 0;
    85  }
    86  };
    87 
    88  //
    89  // Data members
    90  //
    91 
    92  Params params;
    93 
    94  //
    95  // Methods
    96  //
    97 
    99  CUTLASS_DEVICE LinearScaling() { }
    100 
    102  CUTLASS_DEVICE LinearScaling(Params const& _params) : params(_params) {}
    103 
    107  CUTLASS_DEVICE
    108  bool source_required() const {
    109  return !is_zero(params.beta);
    110  }
    111 
    113  template <typename FragmentA_, typename FragmentB_>
    114  CUTLASS_DEVICE void evaluate(FragmentA_ const& accum, FragmentB_& output) {
    115  FragmentMultiplyAdd mad;
    116  mad.multiply(params.alpha, accum, output);
    117 
    118  }
    119 
    121  template <typename ScalarAccum, typename ScalarOutput, int size>
    122  CUTLASS_DEVICE void evaluate(ScalarAccum const *accum, ScalarOutput *output) {
    123  Fragment<ScalarAccum, size> FragAccum;
    124  Fragment<ScalarOutput, size> FragOutput;
    125 #pragma unroll
    126  for (int i = 0; i < size; i++) {
    127  FragAccum[i] = accum[i];
    128  FragOutput[i] = output[i];
    129  }
    130  evaluate(FragAccum, FragOutput);
    131 #pragma unroll
    132  for (int i = 0; i < size; i++) {
    133  output[i] = FragOutput[i];
    134  }
    135  }
    136 
    138  template <typename FragmentA_, typename FragmentB_>
    139  CUTLASS_DEVICE void evaluate(FragmentA_ const& accum, FragmentB_ const& old, FragmentB_& output) {
    140  FragmentMultiplyAdd mad;
    141  FragmentB_ tmp;
    142  mad.multiply(params.beta, old, tmp);
    143  mad.multiply_add(params.alpha, accum, tmp, output);
    144  }
    145 
    147  template <typename ScalarAccum, typename ScalarOutput, int size>
    148  CUTLASS_DEVICE void evaluate(ScalarAccum const *accum, ScalarOutput const *old, ScalarOutput *output) {
    149  Fragment<ScalarAccum, size> FragAccum;
    150  Fragment<ScalarOutput, size> FragOutput;
    151  Fragment<ScalarOutput, size> FragOld;
    152 #pragma unroll
    153  for (int i = 0; i < size; i++) {
    154  FragAccum[i] = accum[i];
    155  FragOutput[i] = output[i];
    156  FragOld[i] = old[i];
    157  }
    158  evaluate(FragAccum, FragOld, FragOutput);
    159 #pragma unroll
    160  for (int i = 0; i < size; i++) {
    161  output[i] = FragOutput[i];
    162  }
    163  }
    164 };
    165 
    167 
    168 } // namespace gemm
    169 } // namespace cutlass
    cutlass::gemm::LinearScaling::Params::initialize
    CUTLASS_HOST_DEVICE int initialize(Scalar _alpha, Scalar _beta)
    Initialize the parameters.
    Definition: linear_scaling.h:73
    +
    cutlass
    Definition: convert.h:33
    +
    cutlass::gemm::LinearScaling::Params::alpha
    Scalar alpha
    The alpha/beta scaling params.
    Definition: linear_scaling.h:62
    +
    cutlass::gemm::LinearScaling::source_required
    CUTLASS_DEVICE bool source_required() const
    Definition: linear_scaling.h:108
    +
    cutlass::gemm::LinearScaling::evaluate
    CUTLASS_DEVICE void evaluate(ScalarAccum const *accum, ScalarOutput *output)
    Evaluate the functor, without using fragment in the API.
    Definition: linear_scaling.h:122
    +
    cutlass::gemm::LinearScaling::evaluate
    CUTLASS_DEVICE void evaluate(FragmentA_ const &accum, FragmentB_ const &old, FragmentB_ &output)
    Evaluate the functor.
    Definition: linear_scaling.h:139
    +
    cutlass::gemm::LinearScaling::evaluate
    CUTLASS_DEVICE void evaluate(FragmentA_ const &accum, FragmentB_ &output)
    Evaluate the functor.
    Definition: linear_scaling.h:114
    +
    cutlass::gemm::LinearScaling::Params::beta
    Scalar beta
    Definition: linear_scaling.h:62
    +
    cutlass::Fragment
    A template defining Fragment Concept.
    Definition: fragment.h:99
    +
    cutlass::gemm::LinearScaling::params
    Params params
    Definition: linear_scaling.h:92
    +
    cutlass::gemm::LinearScaling::ScalarAccum
    FragmentMultiplyAdd_::ScalarAccum ScalarAccum
    Definition: linear_scaling.h:55
    +
    cutlass::gemm::LinearScaling::Params::initialize
    CUTLASS_HOST_DEVICE int initialize(GemmDesc_ const &desc)
    Initialize the parameters.
    Definition: linear_scaling.h:81
    fragment_multiply_add.h
    Defines multiply-add operations on fragments within a thread.
    -
    cutlass::gemm::LinearScaling::FragmentMultiplyAdd
    FragmentMultiplyAdd_ FragmentMultiplyAdd
    Definition: linear_scaling.h:44
    +
    cutlass::gemm::LinearScaling::FragmentMultiplyAdd
    FragmentMultiplyAdd_ FragmentMultiplyAdd
    Definition: linear_scaling.h:57
    +
    cutlass::gemm::LinearScaling::LinearScaling
    CUTLASS_DEVICE LinearScaling()
    Ctor.
    Definition: linear_scaling.h:99
    +
    cutlass::gemm::is_zero
    CUTLASS_DEVICE bool is_zero(T x)
    Definition: linear_scaling.h:39
    CUTLASS_HOST_DEVICE
    #define CUTLASS_HOST_DEVICE
    Definition: cutlass.h:46
    -
    cutlass::gemm::LinearScaling::evaluate
    CUTLASS_DEVICE void evaluate(Fragment_ const &accum, Fragment_ &output)
    Evaluate the functor.
    Definition: linear_scaling.h:65
    -
    cutlass::gemm::LinearScaling::Params
    The parameters.
    Definition: linear_scaling.h:47
    -
    cutlass::gemm::LinearScaling
    Functor to compute linear combination of fragments.
    Definition: linear_scaling.h:40
    -
    cutlass::gemm::LinearScaling::Scalar
    Scalar_ Scalar
    Definition: linear_scaling.h:42
    +
    cutlass::gemm::LinearScaling::LinearScaling
    CUTLASS_DEVICE LinearScaling(Params const &_params)
    Ctor.
    Definition: linear_scaling.h:102
    +
    cutlass::gemm::LinearScaling::Params
    The parameters.
    Definition: linear_scaling.h:60
    +
    cutlass::gemm::LinearScaling
    Functor to compute linear combination of fragments.
    Definition: linear_scaling.h:51
    +
    cutlass::gemm::LinearScaling::Scalar
    Scalar_ Scalar
    Definition: linear_scaling.h:53
    +
    cutlass::gemm::LinearScaling::evaluate
    CUTLASS_DEVICE void evaluate(ScalarAccum const *accum, ScalarOutput const *old, ScalarOutput *output)
    Evaluate the functor, without using fragment in the API.
    Definition: linear_scaling.h:148
    +
    cutlass::gemm::LinearScaling::Params::Params
    CUTLASS_HOST_DEVICE Params(Scalar _alpha=0, Scalar _beta=0)
    Definition: linear_scaling.h:70
    -Generated on Thu May 3 2018 16:36:12 for Cutlass by   +Generated on Wed Sep 19 2018 13:58:50 for Cutlass by   doxygen 1.8.14
    diff --git a/docs/linear__scaling__device__ptr_8h.html b/docs/linear__scaling__device__ptr_8h.html new file mode 100644 index 00000000..ad2add1d --- /dev/null +++ b/docs/linear__scaling__device__ptr_8h.html @@ -0,0 +1,114 @@ + + + + + + + +Cutlass: linear_scaling_device_ptr.h File Reference + + + + + + + + + + +
    +
    + + + + + + +
    +
    Cutlass +
    +
    CUDA Templates for Linear Algebra Subroutines and Solvers
    +
    +
    + + + + + + + + +
    +
    + + +
    + +
    + + +
    +
    +
    +Classes | +Namespaces
    +
    +
    linear_scaling_device_ptr.h File Reference
    +
    +
    + +

    Implements the BLAS linear scaling function alpha*AB + beta*C. +More...

    +
    #include "cutlass/cutlass.h"
    +#include "cutlass/gemm/scalar_or_pointer.h"
    +#include "cutlass/gemm/linear_scaling.h"
    +
    +

    Go to the source code of this file.

    + + + + + + + +

    +Classes

    struct  cutlass::gemm::LinearScalingDevicePtr< Scalar_, FragmentMultiplyAdd_ >
     
    class  cutlass::gemm::LinearScalingDevicePtr< Scalar_, FragmentMultiplyAdd_ >::Params
     The parameters. More...
     
    + + + + + +

    +Namespaces

     cutlass
     
     cutlass::gemm
     
    +
    + +
    +Generated on Wed Sep 19 2018 13:58:50 for Cutlass by   +doxygen + 1.8.14 +
    + + diff --git a/docs/linear__scaling__device__ptr_8h_source.html b/docs/linear__scaling__device__ptr_8h_source.html new file mode 100644 index 00000000..2fae588f --- /dev/null +++ b/docs/linear__scaling__device__ptr_8h_source.html @@ -0,0 +1,109 @@ + + + + + + + +Cutlass: linear_scaling_device_ptr.h Source File + + + + + + + + + + +
    +
    + + + + + + +
    +
    Cutlass +
    +
    CUDA Templates for Linear Algebra Subroutines and Solvers
    +
    +
    + + + + + + + + +
    +
    + + +
    + +
    + + +
    +
    +
    +
    linear_scaling_device_ptr.h
    +
    +
    +Go to the documentation of this file.
    1 /***************************************************************************************************
    2  * Copyright (c) 2017-2018, NVIDIA CORPORATION. All rights reserved.
    3  *
    4  * Redistribution and use in source and binary forms, with or without modification, are permitted
    5  * provided that the following conditions are met:
    6  * * Redistributions of source code must retain the above copyright notice, this list of
    7  * conditions and the following disclaimer.
    8  * * Redistributions in binary form must reproduce the above copyright notice, this list of
    9  * conditions and the following disclaimer in the documentation and/or other materials
    10  * provided with the distribution.
    11  * * Neither the name of the NVIDIA CORPORATION nor the names of its contributors may be used
    12  * to endorse or promote products derived from this software without specific prior written
    13  * permission.
    14  *
    15  * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND ANY EXPRESS OR
    16  * IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND
    17  * FITNESS FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE
    18  * FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING,
    19  * BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS;
    20  * OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT,
    21  * STRICT LIABILITY, OR TOR (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
    22  * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
    23  *
    24  **************************************************************************************************/
    28 #pragma once
    29 
    30 #include "cutlass/cutlass.h"
    31 #include "cutlass/gemm/scalar_or_pointer.h"
    32 #include "cutlass/gemm/linear_scaling.h"
    33 
    34 namespace cutlass {
    35 
    37 
    38 namespace gemm {
    39 
    41 
    45 template <typename Scalar_, typename FragmentMultiplyAdd_ = FragmentMultiplyAdd<Scalar_, Scalar_> >
    46 struct LinearScalingDevicePtr : public LinearScaling<Scalar_, FragmentMultiplyAdd_> {
    47 
    49  typedef LinearScaling<Scalar_, FragmentMultiplyAdd_> Base;
    50 
    51  // The scalar.
    52  typedef typename Base::Scalar Scalar;
    53 
    55  class Params {
    56  private:
    58  detail::ScalarOrPointer<Scalar> alpha_;
    59 
    61  detail::ScalarOrPointer<Scalar> beta_;
    62 
    63  public:
    64  //
    65  // Methods
    66  //
    67 
    68  // Constructor
    69  CUTLASS_HOST_DEVICE
    70  Params() {}
    71 
    72  // Constructor
    73  CUTLASS_HOST_DEVICE
    74  Params(
    75  Scalar alpha,
    76  Scalar beta
    77  ):
    78  alpha_(alpha),
    79  beta_(beta) {}
    80 
    81  // Constructor
    82  CUTLASS_HOST_DEVICE
    83  Params(
    84  Scalar const *alpha_ptr,
    85  Scalar const *beta_ptr
    86  ):
    87  alpha_(alpha_ptr),
    88  beta_(alpha_ptr) {}
    89 
    91  CUTLASS_HOST_DEVICE int initialize(
    92  Scalar alpha,
    93  Scalar beta) {
    94 
    95  alpha_ = alpha;
    96  beta_ = beta;
    97 
    98  return 0;
    99  }
    100 
    102  CUTLASS_HOST_DEVICE int initialize(
    103  Scalar const *alpha,
    104  Scalar const *beta) {
    105 
    106  alpha_ = alpha;
    107  beta_= beta;
    108 
    109  return 0;
    110  }
    111 
    113  template <typename GemmDesc_>
    114  CUTLASS_HOST_DEVICE int initialize(GemmDesc_ const& desc) {
    115 
    116  alpha_ = desc.alpha;
    117  beta_ = desc.beta;
    118 
    119  return 0;
    120  }
    121 
    123  CUTLASS_HOST_DEVICE
    124  Scalar alpha() const {
    125  return alpha_;
    126  }
    127 
    129  CUTLASS_HOST_DEVICE
    130  Scalar beta() const {
    131  return beta_;
    132  }
    133  };
    134 
    135  //
    136  // Methods
    137  //
    138 
    140  CUTLASS_HOST_DEVICE LinearScalingDevicePtr(Params const& _params) {
    141  this->params.alpha = _params.alpha();
    142  this->params.beta = _params.beta();
    143  }
    144 };
    145 
    147 
    148 } // namespace gemm
    149 } // namespace cutlass
    cutlass::gemm::LinearScalingDevicePtr::Params::initialize
    CUTLASS_HOST_DEVICE int initialize(Scalar const *alpha, Scalar const *beta)
    Initialize the parameters.
    Definition: linear_scaling_device_ptr.h:102
    +
    cutlass::gemm::LinearScalingDevicePtr::Params
    The parameters.
    Definition: linear_scaling_device_ptr.h:55
    +
    cutlass
    Definition: convert.h:33
    +
    cutlass::gemm::LinearScalingDevicePtr::Params::Params
    CUTLASS_HOST_DEVICE Params(Scalar const *alpha_ptr, Scalar const *beta_ptr)
    Definition: linear_scaling_device_ptr.h:83
    +
    linear_scaling.h
    Implements the BLAS linear scaling function alpha*AB + beta*C.
    +
    scalar_or_pointer.h
    Implements the BLAS linear scaling function alpha*AB + beta*C.
    +
    cutlass::gemm::LinearScalingDevicePtr::Params::initialize
    CUTLASS_HOST_DEVICE int initialize(Scalar alpha, Scalar beta)
    Initialize the parameters.
    Definition: linear_scaling_device_ptr.h:91
    +
    cutlass::gemm::LinearScaling::params
    Params params
    Definition: linear_scaling.h:92
    +
    cutlass::gemm::LinearScalingDevicePtr::Base
    LinearScaling< Scalar_, FragmentMultiplyAdd_ > Base
    Linear Scaling class used.
    Definition: linear_scaling_device_ptr.h:49
    +
    cutlass::gemm::LinearScalingDevicePtr::Params::Params
    CUTLASS_HOST_DEVICE Params()
    Definition: linear_scaling_device_ptr.h:70
    +
    CUTLASS_HOST_DEVICE
    #define CUTLASS_HOST_DEVICE
    Definition: cutlass.h:46
    +
    cutlass::gemm::LinearScalingDevicePtr::Params::Params
    CUTLASS_HOST_DEVICE Params(Scalar alpha, Scalar beta)
    Definition: linear_scaling_device_ptr.h:74
    +
    cutlass::gemm::LinearScalingDevicePtr::Params::beta
    CUTLASS_HOST_DEVICE Scalar beta() const
    Gets the beta scalar.
    Definition: linear_scaling_device_ptr.h:130
    +
    cutlass::gemm::LinearScalingDevicePtr::LinearScalingDevicePtr
    CUTLASS_HOST_DEVICE LinearScalingDevicePtr(Params const &_params)
    Ctor.
    Definition: linear_scaling_device_ptr.h:140
    +
    cutlass::gemm::LinearScalingDevicePtr::Params::alpha
    CUTLASS_HOST_DEVICE Scalar alpha() const
    Gets the alpha scalar.
    Definition: linear_scaling_device_ptr.h:124
    +
    cutlass::gemm::LinearScalingDevicePtr
    Definition: linear_scaling_device_ptr.h:46
    +
    cutlass::gemm::LinearScaling
    Functor to compute linear combination of fragments.
    Definition: linear_scaling.h:51
    +
    cutlass::gemm::LinearScaling::Scalar
    Scalar_ Scalar
    Definition: linear_scaling.h:53
    +
    cutlass::gemm::LinearScalingDevicePtr::Scalar
    Base::Scalar Scalar
    Definition: linear_scaling_device_ptr.h:52
    +
    cutlass.h
    Basic include for CUTLASS macros.
    +
    cutlass::gemm::LinearScalingDevicePtr::Params::initialize
    CUTLASS_HOST_DEVICE int initialize(GemmDesc_ const &desc)
    Initialize the parameters.
    Definition: linear_scaling_device_ptr.h:114
    + +
    + +
    +Generated on Wed Sep 19 2018 13:58:50 for Cutlass by   +doxygen + 1.8.14 +
    + + diff --git a/docs/load__store_8h.html b/docs/load__store_8h.html index b23ec3cb..30a4e733 100644 --- a/docs/load__store_8h.html +++ b/docs/load__store_8h.html @@ -82,7 +82,7 @@ $(function() {

    Defines abstractions for efficiently loading and storing vectors to memory. More...

    -
    #include <cutlass/vector.h>
    +
    #include "cutlass/vector.h"

    Go to the source code of this file.

    @@ -91,25 +91,43 @@ Classes - + + - + - + + - + - + - + - + - + - + - + + + + + + + + + + + + + + + + +
    struct  cutlass::MemorySpace
     Enum to specify which memory space data resides in. More...
     
    struct  cutlass::Load< Scalar_, Lanes_, Memory_, bool, size_t >
    struct  cutlass::FragmentElementType
     Specifies whether iterator storage fragment consists of Scalar values or WMMA matrix. More...
     
    struct  cutlass::Load< Scalar_, Lanes_, Memory_, true, 4 >
    struct  cutlass::Load< Scalar_, kAccessSize, Memory_, kFragmentElementType, FragmentElement_, kStride, size >
     
    struct  cutlass::Load< Scalar_, Lanes_, Memory_, true, 8 >
    struct  cutlass::Load< Scalar_, kAccessSize, Memory_, FragmentElementType::kScalar, Scalar_, 1, 2 >
     Partial specialization for 16b loads. More...
     
    struct  cutlass::Load< double, 2, Memory_, true, 16 >
    struct  cutlass::Load< Scalar_, kAccessSize, Memory_, FragmentElementType::kScalar, Scalar_, kStride, 4 >
     
    struct  cutlass::Load< Scalar_, Lanes_, Memory_, true, 16 >
    struct  cutlass::Load< Scalar_, kAccessSize, Memory_, FragmentElementType::kScalar, Scalar_, kStride, 8 >
     
    struct  cutlass::Store< Scalar_, Lanes_, Memory_, bool, size_t >
    struct  cutlass::Load< double, 2, Memory_, FragmentElementType::kScalar, double, kStride, 16 >
     
    struct  cutlass::Store< Scalar_, Lanes_, Memory_, true, 4 >
    struct  cutlass::Load< Scalar_, kAccessSize, Memory_, FragmentElementType::kScalar, Scalar_, kStride, 16 >
     
    struct  cutlass::Store< Scalar_, Lanes_, Memory_, true, 8 >
    struct  cutlass::Store< Scalar_, kAccessSize, Memory_, kFragmentElementType, FragmentElement_, kStride, size >
     
    struct  cutlass::Store< double, 2, Memory_, true, 16 >
    struct  cutlass::Store< Scalar_, kAccessSize, Memory_, FragmentElementType::kScalar, Scalar_, 1, 2 >
     
    struct  cutlass::Store< Scalar_, Lanes_, Memory_, true, 16 >
    struct  cutlass::Store< Scalar_, kAccessSize, Memory_, FragmentElementType::kScalar, Scalar_, kStride, 4 >
     
    struct  cutlass::Store< Scalar_, kAccessSize, Memory_, FragmentElementType::kScalar, Scalar_, kStride, 8 >
     
    struct  cutlass::Store< double, 2, Memory_, FragmentElementType::kScalar, double, kStride, 16 >
     
    struct  cutlass::Store< Scalar_, kAccessSize, Memory_, FragmentElementType::kScalar, Scalar_, kStride, 16 >
     
    struct  cutlass::Load< Scalar_, kAccessSize, Memory_, FragmentElementType::kWmmaMatrix, FragmentElement_, kStride, size >
     
    struct  cutlass::Load< Vector< bin1_t, 32 >, kAccessSize, Memory_, FragmentElementType::kWmmaMatrix, FragmentElement_, kStride, size >
     
    struct  cutlass::Load< Vector< int4_t, 8 >, kAccessSize, Memory_, FragmentElementType::kWmmaMatrix, FragmentElement_, kStride, size >
     
    struct  cutlass::Load< Vector< uint4_t, 8 >, kAccessSize, Memory_, FragmentElementType::kWmmaMatrix, FragmentElement_, kStride, size >
     
    struct  cutlass::Store< Scalar_, kAccessSize, Memory_, FragmentElementType::kWmmaMatrix, FragmentElement_, kStride, size >
     
    -Generated on Thu May 3 2018 16:36:12 for Cutlass by   +Generated on Wed Sep 19 2018 13:58:50 for Cutlass by   doxygen 1.8.14
    diff --git a/docs/load__store_8h_source.html b/docs/load__store_8h_source.html index e421cbf2..9fc9c866 100644 --- a/docs/load__store_8h_source.html +++ b/docs/load__store_8h_source.html @@ -76,41 +76,64 @@ $(function() {
    load_store.h
    -Go to the documentation of this file.
    1 /***************************************************************************************************
    2  * Copyright (c) 2017, NVIDIA CORPORATION. All rights reserved.
    3  *
    4  * Redistribution and use in source and binary forms, with or without modification, are permitted
    5  * provided that the following conditions are met:
    6  * * Redistributions of source code must retain the above copyright notice, this list of
    7  * conditions and the following disclaimer.
    8  * * Redistributions in binary form must reproduce the above copyright notice, this list of
    9  * conditions and the following disclaimer in the documentation and/or other materials
    10  * provided with the distribution.
    11  * * Neither the name of the NVIDIA CORPORATION nor the names of its contributors may be used
    12  * to endorse or promote products derived from this software without specific prior written
    13  * permission.
    14  *
    15  * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND ANY EXPRESS OR
    16  * IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND
    17  * FITNESS FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE
    18  * FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING,
    19  * BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS;
    20  * OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT,
    21  * STRICT LIABILITY, OR TOR (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
    22  * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
    23  *
    24  **************************************************************************************************/
    28 #pragma once
    29 
    30 #include <cutlass/vector.h>
    31 
    32 namespace cutlass {
    33 
    35 
    39 struct MemorySpace {
    40  enum Kind {
    41  kGeneric, // Data accessed through pointer dereferencing
    42  kShared, // Data resides in shared memory
    43  kGlobal // Data resides in global memory
    44  };
    45 };
    46 
    48 
    49 template <typename Scalar_,
    50  int Lanes_,
    51  MemorySpace::Kind Memory_,
    52  bool = (Lanes_ > 1),
    53  size_t = (sizeof(Scalar_) * Lanes_)>
    54 struct Load {
    56  typedef typename Vectorize<Scalar_, Lanes_>::Type AccessType;
    57 
    59  static CUTLASS_DEVICE void load(AccessType& dst, Scalar_ const* pointer, int offset) {
    60  dst = reinterpret_cast<AccessType const*>(&pointer[offset])[0];
    61  }
    62 };
    63 
    65 
    66 template <typename Scalar_, int Lanes_, MemorySpace::Kind Memory_>
    67 struct Load<Scalar_, Lanes_, Memory_, true, 4> {
    69  typedef typename Vectorize<Scalar_, Lanes_>::Type AccessType;
    70 
    72  static CUTLASS_DEVICE void load(AccessType& dst, Scalar_ const* pointer, int offset) {
    73  dst.registers[0] = reinterpret_cast<uint32_t const*>(&pointer[offset])[0];
    74  }
    75 };
    76 
    78 
    79 template <typename Scalar_, int Lanes_, MemorySpace::Kind Memory_>
    80 struct Load<Scalar_, Lanes_, Memory_, true, 8> {
    82  typedef typename Vectorize<Scalar_, Lanes_>::Type AccessType;
    83 
    85  static CUTLASS_DEVICE void load(AccessType& dst, Scalar_ const* pointer, int offset) {
    86  uint2 tmp = reinterpret_cast<uint2 const*>(&pointer[offset])[0];
    87  dst.registers[0] = tmp.x;
    88  dst.registers[1] = tmp.y;
    89  }
    90 };
    91 
    93 
    94 template <MemorySpace::Kind Memory_>
    95 struct Load<double, 2, Memory_, true, 16> {
    97  typedef typename Vectorize<double, 2>::Type AccessType;
    98 
    100  static CUTLASS_DEVICE void load(AccessType& dst, double const* pointer, int offset) {
    101  double2 tmp = reinterpret_cast<double2 const*>(&pointer[offset])[0];
    102  dst[0] = tmp.x;
    103  dst[1] = tmp.y;
    104  }
    105 };
    106 
    108 
    109 template <typename Scalar_, int Lanes_, MemorySpace::Kind Memory_>
    110 struct Load<Scalar_, Lanes_, Memory_, true, 16> {
    112  typedef typename Vectorize<Scalar_, Lanes_>::Type AccessType;
    113 
    115  static CUTLASS_DEVICE void load(AccessType& dst, Scalar_ const* pointer, int offset) {
    116  uint4 tmp = reinterpret_cast<uint4 const*>(&pointer[offset])[0];
    117  dst.registers[0] = tmp.x;
    118  dst.registers[1] = tmp.y;
    119  dst.registers[2] = tmp.z;
    120  dst.registers[3] = tmp.w;
    121  }
    122 };
    123 
    125 
    126 template <typename Scalar_,
    127  int Lanes_,
    128  MemorySpace::Kind Memory_,
    129  bool = (Lanes_ > 1),
    130  size_t = (sizeof(Scalar_) * Lanes_)>
    131 struct Store {
    133  typedef typename Vectorize<Scalar_, Lanes_>::Type AccessType;
    134 
    136  static CUTLASS_DEVICE void store(AccessType const& src, Scalar_* pointer, int offset) {
    137  pointer[offset] = src;
    138  }
    139 };
    140 
    142 
    143 template <typename Scalar_, int Lanes_, MemorySpace::Kind Memory_>
    144 struct Store<Scalar_, Lanes_, Memory_, true, 4> {
    146  typedef typename Vectorize<Scalar_, Lanes_>::Type AccessType;
    147 
    149  static CUTLASS_DEVICE void store(AccessType const& src, Scalar_* pointer, int offset) {
    150  uint32_t* addr = reinterpret_cast<uint32_t*>(&pointer[offset]);
    151  addr[0] = src.registers[0];
    152  }
    153 };
    154 
    156 
    157 template <typename Scalar_, int Lanes_, MemorySpace::Kind Memory_>
    158 struct Store<Scalar_, Lanes_, Memory_, true, 8> {
    160  typedef typename Vectorize<Scalar_, Lanes_>::Type AccessType;
    161 
    163  static CUTLASS_DEVICE void store(AccessType const& src, Scalar_* pointer, int offset) {
    164  uint2* addr = reinterpret_cast<uint2*>(&pointer[offset]);
    165  addr[0] = make_uint2(src.registers[0], src.registers[1]);
    166  }
    167 };
    168 
    170 
    171 template <MemorySpace::Kind Memory_>
    172 struct Store<double, 2, Memory_, true, 16> {
    174  typedef typename Vectorize<double, 2>::Type AccessType;
    175 
    177  static CUTLASS_DEVICE void store(AccessType const& src, double* pointer, int offset) {
    178  double2* addr = reinterpret_cast<double2*>(&pointer[offset]);
    179  addr[0] = make_double2(src[0], src[1]);
    180  }
    181 };
    182 
    184 
    185 template <typename Scalar_, int Lanes_, MemorySpace::Kind Memory_>
    186 struct Store<Scalar_, Lanes_, Memory_, true, 16> {
    188  typedef typename Vectorize<Scalar_, Lanes_>::Type AccessType;
    189 
    191  static CUTLASS_DEVICE void store(AccessType const& src, Scalar_* pointer, int offset) {
    192  uint4* addr = reinterpret_cast<uint4*>(&pointer[offset]);
    193  addr[0] = make_uint4(src.registers[0], src.registers[1], src.registers[2], src.registers[3]);
    194  }
    195 };
    196 
    198 
    199 } // namespace cutlass
    cutlass::Store< Scalar_, Lanes_, Memory_, true, 16 >::AccessType
    Vectorize< Scalar_, Lanes_ >::Type AccessType
    The output type.
    Definition: load_store.h:188
    -
    cutlass::MemorySpace::kShared
    Definition: load_store.h:42
    +Go to the documentation of this file.
    1 /***************************************************************************************************
    2  * Copyright (c) 2017, NVIDIA CORPORATION. All rights reserved.
    3  *
    4  * Redistribution and use in source and binary forms, with or without modification, are permitted
    5  * provided that the following conditions are met:
    6  * * Redistributions of source code must retain the above copyright notice, this list of
    7  * conditions and the following disclaimer.
    8  * * Redistributions in binary form must reproduce the above copyright notice, this list of
    9  * conditions and the following disclaimer in the documentation and/or other materials
    10  * provided with the distribution.
    11  * * Neither the name of the NVIDIA CORPORATION nor the names of its contributors may be used
    12  * to endorse or promote products derived from this software without specific prior written
    13  * permission.
    14  *
    15  * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND ANY EXPRESS OR
    16  * IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND
    17  * FITNESS FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE
    18  * FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING,
    19  * BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS;
    20  * OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT,
    21  * STRICT LIABILITY, OR TOR (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
    22  * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
    23  *
    24  **************************************************************************************************/
    28 #pragma once
    29 
    30 #include "cutlass/vector.h"
    31 namespace cutlass {
    32 
    34 
    38 struct MemorySpace {
    39  enum Kind {
    40  kGeneric, // Data accessed through pointer dereferencing
    41  kShared, // Data resides in shared memory
    42  kGlobal // Data resides in global memory
    43  };
    44 };
    45 
    47 struct FragmentElementType {
    48  enum Kind { kScalar, kWmmaMatrix };
    49 };
    50 
    52 
    53 template <typename Scalar_,
    54  int kAccessSize,
    55  MemorySpace::Kind Memory_,
    56  FragmentElementType::Kind kFragmentElementType = FragmentElementType::kScalar,
    57  typename FragmentElement_ = Scalar_,
    58  int kStride = 1,
    59  size_t size = (sizeof(Scalar_) * kAccessSize)>
    60 struct Load {
    62  typedef typename Vectorize<Scalar_, kAccessSize>::Type AccessType;
    63 
    65  static CUTLASS_HOST_DEVICE void load(AccessType& dst, Scalar_ const* pointer, int offset) {
    66  dst = *reinterpret_cast<AccessType const*>(pointer + offset);
    67  }
    68 
    69 };
    70 
    72 
    74 template <typename Scalar_, int kAccessSize, MemorySpace::Kind Memory_>
    75 struct Load<Scalar_, kAccessSize, Memory_, FragmentElementType::kScalar, Scalar_, 1, 2> {
    77  typedef typename Vectorize<Scalar_, kAccessSize>::Type AccessType;
    78 
    80  static CUTLASS_HOST_DEVICE void load(AccessType& dst, Scalar_ const* pointer, int offset) {
    81  reinterpret_cast<uint16_t&>(dst) = reinterpret_cast<uint16_t const*>(&pointer[offset])[0];
    82  }
    83 };
    84 
    86 
    87 template <typename Scalar_, int kAccessSize, MemorySpace::Kind Memory_, int kStride>
    88 struct Load<Scalar_, kAccessSize, Memory_, FragmentElementType::kScalar, Scalar_, kStride, 4> {
    90  typedef typename Vectorize<Scalar_, kAccessSize>::Type AccessType;
    91 
    93  static CUTLASS_HOST_DEVICE void load(AccessType& dst, Scalar_ const* pointer, int offset) {
    94  dst.registers[0] = reinterpret_cast<uint32_t const*>(&pointer[offset])[0];
    95  }
    96 
    97 };
    98 
    100 
    101 template <typename Scalar_, int kAccessSize, MemorySpace::Kind Memory_, int kStride>
    102 struct Load<Scalar_, kAccessSize, Memory_, FragmentElementType::kScalar, Scalar_, kStride, 8> {
    104  typedef typename Vectorize<Scalar_, kAccessSize>::Type AccessType;
    105 
    107  static CUTLASS_HOST_DEVICE void load(AccessType& dst, Scalar_ const* pointer, int offset) {
    108  uint2 tmp = reinterpret_cast<uint2 const*>(&pointer[offset])[0];
    109  dst.registers[0] = tmp.x;
    110  dst.registers[1] = tmp.y;
    111  }
    112 };
    113 
    115 
    116 template <MemorySpace::Kind Memory_, int kStride>
    117 struct Load<double, 2, Memory_, FragmentElementType::kScalar, double, kStride, 16> {
    119  typedef typename Vectorize<double, 2>::Type AccessType;
    120 
    122  static CUTLASS_HOST_DEVICE void load(AccessType& dst, double const* pointer, int offset) {
    123  double2 tmp = reinterpret_cast<double2 const*>(&pointer[offset])[0];
    124  dst[0] = tmp.x;
    125  dst[1] = tmp.y;
    126  }
    127 };
    128 
    130 
    131 #if defined(__CUDACC_VERSION_MAJOR) && __CUDACC_VERSION_MAJOR < 10
    132 // WAR bug in NVCC where the upper and lower half of the register end up being the same
    133 template <MemorySpace::Kind Memory_, int kStride>
    134 struct Load<half, 8, Memory_, FragmentElementType::kScalar, half, kStride, 16> {
    136  typedef typename Vectorize<half, 8>::Type AccessType;
    137 
    139  static CUTLASS_HOST_DEVICE void load(AccessType& dst, half const* pointer, int offset) {
    140  int2 tmp = reinterpret_cast<int2 const*>(&pointer[offset])[0];
    141  dst.registers[0] = tmp.x;
    142  dst.registers[1] = tmp.y;
    143 
    144  tmp = reinterpret_cast<int2 const*>(&pointer[offset + 4])[0];
    145  dst.registers[2] = tmp.x;
    146  dst.registers[3] = tmp.y;
    147  }
    148 };
    149 
    150 #endif
    151 
    153 
    154 template <typename Scalar_, int kAccessSize, MemorySpace::Kind Memory_, int kStride>
    155 struct Load<Scalar_, kAccessSize, Memory_, FragmentElementType::kScalar, Scalar_, kStride, 16> {
    157  typedef typename Vectorize<Scalar_, kAccessSize>::Type AccessType;
    158 
    160  static CUTLASS_HOST_DEVICE void load(AccessType& dst, Scalar_ const* pointer, int offset) {
    161  uint4 tmp = reinterpret_cast<uint4 const*>(&pointer[offset])[0];
    162  dst.registers[0] = tmp.x;
    163  dst.registers[1] = tmp.y;
    164  dst.registers[2] = tmp.z;
    165  dst.registers[3] = tmp.w;
    166  }
    167 };
    168 
    170 
    171 template <typename Scalar_,
    172  int kAccessSize,
    173  MemorySpace::Kind Memory_,
    174  FragmentElementType::Kind kFragmentElementType = FragmentElementType::kScalar,
    175  typename FragmentElement_ = Scalar_,
    176  int kStride = 1,
    177  size_t size = (sizeof(Scalar_) * kAccessSize)>
    178 struct Store {
    180  typedef typename Vectorize<FragmentElement_, kAccessSize>::Type AccessType;
    181 
    183  static CUTLASS_HOST_DEVICE void store(AccessType const& src, Scalar_* pointer, int offset) {
    184  pointer[offset] = *reinterpret_cast<Scalar_ const*>(&src);
    185  }
    186 };
    187 
    189 
    190 template <typename Scalar_, int kAccessSize, MemorySpace::Kind Memory_>
    191 struct Store<Scalar_, kAccessSize, Memory_, FragmentElementType::kScalar, Scalar_, 1, 2> {
    193  typedef typename Vectorize<Scalar_, kAccessSize>::Type AccessType;
    194 
    196  static CUTLASS_HOST_DEVICE void store(AccessType const& src, Scalar_* pointer, int offset) {
    197  uint16_t* addr = reinterpret_cast<uint16_t*>(&pointer[offset]);
    198  addr[0] = reinterpret_cast<uint16_t const&>(src);
    199  }
    200 };
    201 
    203 
    204 template <typename Scalar_, int kAccessSize, MemorySpace::Kind Memory_, int kStride>
    205 struct Store<Scalar_, kAccessSize, Memory_, FragmentElementType::kScalar, Scalar_, kStride, 4> {
    207  typedef typename Vectorize<Scalar_, kAccessSize>::Type AccessType;
    208 
    210  static CUTLASS_HOST_DEVICE void store(AccessType const& src, Scalar_* pointer, int offset) {
    211  uint32_t* addr = reinterpret_cast<uint32_t*>(&pointer[offset]);
    212  addr[0] = src.registers[0];
    213  }
    214 };
    215 
    217 
    218 template <typename Scalar_, int kAccessSize, MemorySpace::Kind Memory_, int kStride>
    219 struct Store<Scalar_, kAccessSize, Memory_, FragmentElementType::kScalar, Scalar_, kStride, 8> {
    221  typedef typename Vectorize<Scalar_, kAccessSize>::Type AccessType;
    222 
    224  static CUTLASS_HOST_DEVICE void store(AccessType const& src, Scalar_* pointer, int offset) {
    225  uint2* addr = reinterpret_cast<uint2*>(&pointer[offset]);
    226  addr[0] = make_uint2(src.registers[0], src.registers[1]);
    227  }
    228 };
    229 
    231 
    232 template <MemorySpace::Kind Memory_, int kStride>
    233 struct Store<double, 2, Memory_, FragmentElementType::kScalar, double, kStride, 16> {
    235  typedef typename Vectorize<double, 2>::Type AccessType;
    236 
    238  static CUTLASS_HOST_DEVICE void store(AccessType const& src, double* pointer, int offset) {
    239  double2* addr = reinterpret_cast<double2*>(&pointer[offset]);
    240  addr[0] = make_double2(src[0], src[1]);
    241  }
    242 };
    243 
    245 
    246 template <typename Scalar_, int kAccessSize, MemorySpace::Kind Memory_, int kStride>
    247 struct Store<Scalar_, kAccessSize, Memory_, FragmentElementType::kScalar, Scalar_, kStride, 16> {
    249  typedef typename Vectorize<Scalar_, kAccessSize>::Type AccessType;
    250 
    252  static CUTLASS_HOST_DEVICE void store(AccessType const& src, Scalar_* pointer, int offset) {
    253  uint4* addr = reinterpret_cast<uint4*>(&pointer[offset]);
    254  addr[0] = make_uint4(src.registers[0], src.registers[1], src.registers[2], src.registers[3]);
    255  }
    256 };
    257 
    259 
    260 template <typename Scalar_,
    261  int kAccessSize,
    262  MemorySpace::Kind Memory_,
    263  typename FragmentElement_,
    264  int kStride,
    265  size_t size>
    266 struct Load<Scalar_,
    267  kAccessSize,
    268  Memory_,
    269  FragmentElementType::kWmmaMatrix,
    270  FragmentElement_,
    271  kStride,
    272  size> {
    274  typedef FragmentElement_ AccessType;
    275 
    277  static CUTLASS_HOST_DEVICE void load(AccessType& value, Scalar_ const* pointer, int offset) {
    278  value.load(&pointer[offset], kStride);
    279  }
    280 };
    281 
    283 
    284 template <int kAccessSize,
    285  MemorySpace::Kind Memory_,
    286  typename FragmentElement_,
    287  int kStride,
    288  size_t size>
    289 struct Load<Vector<bin1_t, 32>,
    290  kAccessSize,
    291  Memory_,
    292  FragmentElementType::kWmmaMatrix,
    293  FragmentElement_,
    294  kStride,
    295  size> {
    297  typedef FragmentElement_ AccessType;
    298 
    300  static CUTLASS_HOST_DEVICE void load(AccessType& value, Vector<bin1_t, 32> const* pointer,
    301  int offset) {
    302  value.load(&pointer[offset], kStride * 32);
    303  }
    304 };
    305 
    307 
    308 template <int kAccessSize,
    309  MemorySpace::Kind Memory_,
    310  typename FragmentElement_,
    311  int kStride,
    312  size_t size>
    313 struct Load<Vector<int4_t, 8>,
    314  kAccessSize,
    315  Memory_,
    316  FragmentElementType::kWmmaMatrix,
    317  FragmentElement_,
    318  kStride,
    319  size> {
    321  typedef FragmentElement_ AccessType;
    322 
    324  static CUTLASS_HOST_DEVICE void load(AccessType& value, Vector<int4_t, 8> const* pointer,
    325  int offset) {
    326  value.load(&pointer[offset], kStride * 8);
    327  }
    328 };
    329 
    331 
    332 template <int kAccessSize,
    333  MemorySpace::Kind Memory_,
    334  typename FragmentElement_,
    335  int kStride,
    336  size_t size>
    337 struct Load<Vector<uint4_t, 8>,
    338  kAccessSize,
    339  Memory_,
    340  FragmentElementType::kWmmaMatrix,
    341  FragmentElement_,
    342  kStride,
    343  size> {
    345  typedef FragmentElement_ AccessType;
    346 
    348  static CUTLASS_HOST_DEVICE void load(AccessType& value, Vector<uint4_t, 8> const* pointer,
    349  int offset) {
    350  value.load(&pointer[offset], kStride * 8);
    351  }
    352 };
    353 
    355 template <typename Scalar_,
    356  int kAccessSize,
    357  MemorySpace::Kind Memory_,
    358  typename FragmentElement_,
    359  int kStride,
    360  size_t size>
    361 struct Store<Scalar_,
    362  kAccessSize,
    363  Memory_,
    364  FragmentElementType::kWmmaMatrix,
    365  FragmentElement_,
    366  kStride,
    367  size> {
    369  typedef FragmentElement_ AccessType;
    370 
    372  static CUTLASS_HOST_DEVICE void store(AccessType const& value, Scalar_* pointer, int offset) {
    373  value.store(&pointer[offset], kStride);
    374  }
    375 };
    376 
    378 
    379 } // namespace cutlass
    cutlass::Load< Vector< bin1_t, 32 >, kAccessSize, Memory_, FragmentElementType::kWmmaMatrix, FragmentElement_, kStride, size >::load
    static CUTLASS_HOST_DEVICE void load(AccessType &value, Vector< bin1_t, 32 > const *pointer, int offset)
    The load function.
    Definition: load_store.h:300
    +
    cutlass::Load< Scalar_, kAccessSize, Memory_, FragmentElementType::kScalar, Scalar_, kStride, 16 >::AccessType
    Vectorize< Scalar_, kAccessSize >::Type AccessType
    The output type.
    Definition: load_store.h:157
    +
    cutlass::Load< Scalar_, kAccessSize, Memory_, FragmentElementType::kScalar, Scalar_, 1, 2 >::AccessType
    Vectorize< Scalar_, kAccessSize >::Type AccessType
    The output type.
    Definition: load_store.h:77
    +
    cutlass::Store< double, 2, Memory_, FragmentElementType::kScalar, double, kStride, 16 >::store
    static CUTLASS_HOST_DEVICE void store(AccessType const &src, double *pointer, int offset)
    The store function.
    Definition: load_store.h:238
    +
    cutlass::Load< Vector< int4_t, 8 >, kAccessSize, Memory_, FragmentElementType::kWmmaMatrix, FragmentElement_, kStride, size >::load
    static CUTLASS_HOST_DEVICE void load(AccessType &value, Vector< int4_t, 8 > const *pointer, int offset)
    The load function.
    Definition: load_store.h:324
    +
    cutlass::MemorySpace::kShared
    Definition: load_store.h:41
    cutlass
    Definition: convert.h:33
    -
    cutlass::Store< Scalar_, Lanes_, Memory_, true, 8 >::store
    static CUTLASS_DEVICE void store(AccessType const &src, Scalar_ *pointer, int offset)
    The store function.
    Definition: load_store.h:163
    -
    cutlass::MemorySpace
    Enum to specify which memory space data resides in.
    Definition: load_store.h:39
    -
    cutlass::MemorySpace::kGlobal
    Definition: load_store.h:43
    -
    cutlass::Load::load
    static CUTLASS_DEVICE void load(AccessType &dst, Scalar_ const *pointer, int offset)
    The load function.
    Definition: load_store.h:59
    -
    cutlass::Load< Scalar_, Lanes_, Memory_, true, 16 >::AccessType
    Vectorize< Scalar_, Lanes_ >::Type AccessType
    The output type.
    Definition: load_store.h:112
    -
    cutlass::Store< Scalar_, Lanes_, Memory_, true, 4 >::AccessType
    Vectorize< Scalar_, Lanes_ >::Type AccessType
    The output type.
    Definition: load_store.h:146
    -
    cutlass::MemorySpace::Kind
    Kind
    Definition: load_store.h:40
    -
    cutlass::Store
    Definition: load_store.h:131
    -
    cutlass::Store::store
    static CUTLASS_DEVICE void store(AccessType const &src, Scalar_ *pointer, int offset)
    The store function.
    Definition: load_store.h:136
    -
    cutlass::Vector::registers
    uint32_t registers[kRegisters]
    The data in registers.
    Definition: vector.h:80
    -
    cutlass::Store< double, 2, Memory_, true, 16 >::AccessType
    Vectorize< double, 2 >::Type AccessType
    The output type.
    Definition: load_store.h:174
    -
    cutlass::MemorySpace::kGeneric
    Definition: load_store.h:41
    -
    cutlass::Load< Scalar_, Lanes_, Memory_, true, 4 >::load
    static CUTLASS_DEVICE void load(AccessType &dst, Scalar_ const *pointer, int offset)
    The store function.
    Definition: load_store.h:72
    -
    cutlass::Store::AccessType
    Vectorize< Scalar_, Lanes_ >::Type AccessType
    The output type.
    Definition: load_store.h:133
    -
    cutlass::Vector
    Definition: vector.h:61
    -
    cutlass::Load< Scalar_, Lanes_, Memory_, true, 8 >::load
    static CUTLASS_DEVICE void load(AccessType &dst, Scalar_ const *pointer, int offset)
    The store function.
    Definition: load_store.h:85
    -
    cutlass::Load
    Definition: load_store.h:54
    -
    cutlass::Load< Scalar_, Lanes_, Memory_, true, 8 >::AccessType
    Vectorize< Scalar_, Lanes_ >::Type AccessType
    The output type.
    Definition: load_store.h:82
    +
    cutlass::Load< Vector< int4_t, 8 >, kAccessSize, Memory_, FragmentElementType::kWmmaMatrix, FragmentElement_, kStride, size >::AccessType
    FragmentElement_ AccessType
    The output type.
    Definition: load_store.h:321
    +
    cutlass::bin1_t
    Definition: numeric_types.h:39
    +
    cutlass::MemorySpace
    Enum to specify which memory space data resides in.
    Definition: load_store.h:38
    +
    cutlass::Store< Scalar_, kAccessSize, Memory_, FragmentElementType::kScalar, Scalar_, 1, 2 >::store
    static CUTLASS_HOST_DEVICE void store(AccessType const &src, Scalar_ *pointer, int offset)
    The store function.
    Definition: load_store.h:196
    +
    cutlass::Store< Scalar_, kAccessSize, Memory_, FragmentElementType::kScalar, Scalar_, kStride, 16 >::store
    static CUTLASS_HOST_DEVICE void store(AccessType const &src, Scalar_ *pointer, int offset)
    The store function.
    Definition: load_store.h:252
    +
    cutlass::Load< Vector< bin1_t, 32 >, kAccessSize, Memory_, FragmentElementType::kWmmaMatrix, FragmentElement_, kStride, size >::AccessType
    FragmentElement_ AccessType
    The output type.
    Definition: load_store.h:297
    +
    cutlass::FragmentElementType
    Specifies whether iterator storage fragment consists of Scalar values or WMMA matrix.
    Definition: load_store.h:47
    +
    cutlass::MemorySpace::kGlobal
    Definition: load_store.h:42
    + +
    cutlass::Load< Vector< uint4_t, 8 >, kAccessSize, Memory_, FragmentElementType::kWmmaMatrix, FragmentElement_, kStride, size >::AccessType
    FragmentElement_ AccessType
    The output type.
    Definition: load_store.h:345
    +
    cutlass::Load< double, 2, Memory_, FragmentElementType::kScalar, double, kStride, 16 >::AccessType
    Vectorize< double, 2 >::Type AccessType
    The output type.
    Definition: load_store.h:119
    +
    cutlass::Store::AccessType
    Vectorize< FragmentElement_, kAccessSize >::Type AccessType
    The output type.
    Definition: load_store.h:180
    +
    cutlass::MemorySpace::Kind
    Kind
    Definition: load_store.h:39
    +
    cutlass::Store
    Definition: load_store.h:178
    +
    cutlass::Load< Scalar_, kAccessSize, Memory_, FragmentElementType::kScalar, Scalar_, kStride, 16 >::load
    static CUTLASS_HOST_DEVICE void load(AccessType &dst, Scalar_ const *pointer, int offset)
    The load function.
    Definition: load_store.h:160
    +
    cutlass::Vector::registers
    uint32_t registers[kRegisters]
    The data in registers.
    Definition: vector.h:81
    +
    cutlass::Load< Scalar_, kAccessSize, Memory_, FragmentElementType::kWmmaMatrix, FragmentElement_, kStride, size >::load
    static CUTLASS_HOST_DEVICE void load(AccessType &value, Scalar_ const *pointer, int offset)
    The load function.
    Definition: load_store.h:277
    +
    cutlass::Store< Scalar_, kAccessSize, Memory_, FragmentElementType::kScalar, Scalar_, 1, 2 >::AccessType
    Vectorize< Scalar_, kAccessSize >::Type AccessType
    The output type.
    Definition: load_store.h:193
    +
    cutlass::Load< Scalar_, kAccessSize, Memory_, FragmentElementType::kScalar, Scalar_, kStride, 8 >::AccessType
    Vectorize< Scalar_, kAccessSize >::Type AccessType
    The output type.
    Definition: load_store.h:104
    +
    cutlass::FragmentElementType::Kind
    Kind
    Definition: load_store.h:48
    +
    cutlass::MemorySpace::kGeneric
    Definition: load_store.h:40
    +
    CUTLASS_HOST_DEVICE
    #define CUTLASS_HOST_DEVICE
    Definition: cutlass.h:46
    +
    cutlass::Load::AccessType
    Vectorize< Scalar_, kAccessSize >::Type AccessType
    The output type.
    Definition: load_store.h:62
    +
    cutlass::Load< Scalar_, kAccessSize, Memory_, FragmentElementType::kScalar, Scalar_, kStride, 8 >::load
    static CUTLASS_HOST_DEVICE void load(AccessType &dst, Scalar_ const *pointer, int offset)
    The load function.
    Definition: load_store.h:107
    +
    cutlass::Vector
    Definition: vector.h:62
    +
    cutlass::Load
    Definition: load_store.h:60
    +
    cutlass::Store< Scalar_, kAccessSize, Memory_, FragmentElementType::kWmmaMatrix, FragmentElement_, kStride, size >::AccessType
    FragmentElement_ AccessType
    The input type.
    Definition: load_store.h:369
    +
    cutlass::Load< Scalar_, kAccessSize, Memory_, FragmentElementType::kScalar, Scalar_, kStride, 4 >::load
    static CUTLASS_HOST_DEVICE void load(AccessType &dst, Scalar_ const *pointer, int offset)
    The load function.
    Definition: load_store.h:93
    +
    cutlass::FragmentElementType::kScalar
    Definition: load_store.h:48
    +
    cutlass::Vectorize::Type
    Vector< Element_, kLanes_ > Type
    Definition: vector.h:271
    vector.h
    Defines a 1D vector of elements held in the registers of each thread.
    -
    cutlass::Store< Scalar_, Lanes_, Memory_, true, 8 >::AccessType
    Vectorize< Scalar_, Lanes_ >::Type AccessType
    The output type.
    Definition: load_store.h:160
    -
    cutlass::Load< Scalar_, Lanes_, Memory_, true, 16 >::load
    static CUTLASS_DEVICE void load(AccessType &dst, Scalar_ const *pointer, int offset)
    The store function.
    Definition: load_store.h:115
    -
    cutlass::Load< Scalar_, Lanes_, Memory_, true, 4 >::AccessType
    Vectorize< Scalar_, Lanes_ >::Type AccessType
    The output type.
    Definition: load_store.h:69
    -
    cutlass::Load< double, 2, Memory_, true, 16 >::load
    static CUTLASS_DEVICE void load(AccessType &dst, double const *pointer, int offset)
    The store function.
    Definition: load_store.h:100
    -
    cutlass::Load< double, 2, Memory_, true, 16 >::AccessType
    Vectorize< double, 2 >::Type AccessType
    The output type.
    Definition: load_store.h:97
    -
    cutlass::Load::AccessType
    Vectorize< Scalar_, Lanes_ >::Type AccessType
    The output type.
    Definition: load_store.h:56
    -
    cutlass::Store< Scalar_, Lanes_, Memory_, true, 16 >::store
    static CUTLASS_DEVICE void store(AccessType const &src, Scalar_ *pointer, int offset)
    The store function.
    Definition: load_store.h:191
    -
    cutlass::Store< Scalar_, Lanes_, Memory_, true, 4 >::store
    static CUTLASS_DEVICE void store(AccessType const &src, Scalar_ *pointer, int offset)
    The store function.
    Definition: load_store.h:149
    -
    cutlass::Store< double, 2, Memory_, true, 16 >::store
    static CUTLASS_DEVICE void store(AccessType const &src, double *pointer, int offset)
    The store function.
    Definition: load_store.h:177
    +
    cutlass::Store< Scalar_, kAccessSize, Memory_, FragmentElementType::kScalar, Scalar_, kStride, 16 >::AccessType
    Vectorize< Scalar_, kAccessSize >::Type AccessType
    The output type.
    Definition: load_store.h:249
    +
    cutlass::Load< Vector< uint4_t, 8 >, kAccessSize, Memory_, FragmentElementType::kWmmaMatrix, FragmentElement_, kStride, size >::load
    static CUTLASS_HOST_DEVICE void load(AccessType &value, Vector< uint4_t, 8 > const *pointer, int offset)
    The load function.
    Definition: load_store.h:348
    +
    cutlass::uint4_t
    Definition: numeric_types.h:43
    +
    cutlass::Load< Scalar_, kAccessSize, Memory_, FragmentElementType::kScalar, Scalar_, kStride, 4 >::AccessType
    Vectorize< Scalar_, kAccessSize >::Type AccessType
    The output type.
    Definition: load_store.h:90
    +
    cutlass::Store::store
    static CUTLASS_HOST_DEVICE void store(AccessType const &src, Scalar_ *pointer, int offset)
    The store function.
    Definition: load_store.h:183
    +
    cutlass::Store< Scalar_, kAccessSize, Memory_, FragmentElementType::kScalar, Scalar_, kStride, 8 >::store
    static CUTLASS_HOST_DEVICE void store(AccessType const &src, Scalar_ *pointer, int offset)
    The store function.
    Definition: load_store.h:224
    +
    cutlass::Store< Scalar_, kAccessSize, Memory_, FragmentElementType::kScalar, Scalar_, kStride, 8 >::AccessType
    Vectorize< Scalar_, kAccessSize >::Type AccessType
    The output type.
    Definition: load_store.h:221
    +
    cutlass::Store< Scalar_, kAccessSize, Memory_, FragmentElementType::kScalar, Scalar_, kStride, 4 >::AccessType
    Vectorize< Scalar_, kAccessSize >::Type AccessType
    The output type.
    Definition: load_store.h:207
    +
    cutlass::Store< Scalar_, kAccessSize, Memory_, FragmentElementType::kWmmaMatrix, FragmentElement_, kStride, size >::store
    static CUTLASS_HOST_DEVICE void store(AccessType const &value, Scalar_ *pointer, int offset)
    The store function.
    Definition: load_store.h:372
    +
    cutlass::int4_t
    Definition: numeric_types.h:41
    +
    cutlass::Load< Scalar_, kAccessSize, Memory_, FragmentElementType::kWmmaMatrix, FragmentElement_, kStride, size >::AccessType
    FragmentElement_ AccessType
    The output type.
    Definition: load_store.h:274
    +
    cutlass::Store< Scalar_, kAccessSize, Memory_, FragmentElementType::kScalar, Scalar_, kStride, 4 >::store
    static CUTLASS_HOST_DEVICE void store(AccessType const &src, Scalar_ *pointer, int offset)
    The store function.
    Definition: load_store.h:210
    +
    cutlass::Store< double, 2, Memory_, FragmentElementType::kScalar, double, kStride, 16 >::AccessType
    Vectorize< double, 2 >::Type AccessType
    The output type.
    Definition: load_store.h:235
    +
    cutlass::Load::load
    static CUTLASS_HOST_DEVICE void load(AccessType &dst, Scalar_ const *pointer, int offset)
    The load function.
    Definition: load_store.h:65
    +
    cutlass::Load< Scalar_, kAccessSize, Memory_, FragmentElementType::kScalar, Scalar_, 1, 2 >::load
    static CUTLASS_HOST_DEVICE void load(AccessType &dst, Scalar_ const *pointer, int offset)
    The load function.
    Definition: load_store.h:80
    +
    cutlass::Load< double, 2, Memory_, FragmentElementType::kScalar, double, kStride, 16 >::load
    static CUTLASS_HOST_DEVICE void load(AccessType &dst, double const *pointer, int offset)
    The load function.
    Definition: load_store.h:122
    -Generated on Thu May 3 2018 16:36:12 for Cutlass by   +Generated on Wed Sep 19 2018 13:58:50 for Cutlass by   doxygen 1.8.14
    diff --git a/docs/matrix__traits_8h.html b/docs/matrix__traits_8h.html index f83c89f0..097d5f4a 100644 --- a/docs/matrix__traits_8h.html +++ b/docs/matrix__traits_8h.html @@ -82,17 +82,39 @@ $(function() {

    Defines properties of matrices used to denote layout and operands to GEMM kernels. More...

    - +
    #include "cutlass/coord.h"
    +

    Go to the source code of this file.

    @@ -120,7 +138,7 @@ Namespaces

    + + - + + + + + + + + + + + + + + + + + + + +

    Classes

    struct  cutlass::MatrixCoord
     
    struct  cutlass::MatrixLayout
     Describes layouts of matrices. More...
     Defines data layouts of various matrix formats usable by TensorRef and other classes. More...
     
    struct  cutlass::MatrixLayout::RowMajor
     Mapping function for row-major matrices. More...
     
    struct  cutlass::MatrixLayout::ColumnMajor
     Mapping function for column-major matrices. More...
     
    struct  cutlass::MatrixLayout::RowMajorInterleaved< Interleave >
     
    struct  cutlass::MatrixLayout::ColumnMajorInterleaved< Interleave >
     
    struct  cutlass::MatrixLayout::ContiguousLayout
     
    struct  cutlass::MatrixLayout::ColumnMajorBlockLinear< BlockRows, BlockColumns >
     
    struct  cutlass::MatrixLayout::RowMajorBlockLinear< BlockRows, BlockColumns >
     
    struct  cutlass::GemmOperand
     Gemm operand - D = A * B + C. More...
     
    struct  cutlass::MatrixTransform
     Transformation applied to matrix operands. More...
     
    @@ -102,7 +124,7 @@ Namespaces
    -Generated on Thu May 3 2018 16:36:12 for Cutlass by   +Generated on Wed Sep 19 2018 13:58:50 for Cutlass by   doxygen 1.8.14
    diff --git a/docs/matrix__traits_8h_source.html b/docs/matrix__traits_8h_source.html index 9f8de2dc..2e78c5a2 100644 --- a/docs/matrix__traits_8h_source.html +++ b/docs/matrix__traits_8h_source.html @@ -76,21 +76,88 @@ $(function() {
    matrix_traits.h
    -Go to the documentation of this file.
    1 /***************************************************************************************************
    2  * Copyright (c) 2017-2018, NVIDIA CORPORATION. All rights reserved.
    3  *
    4  * Redistribution and use in source and binary forms, with or without modification, are permitted
    5  * provided that the following conditions are met:
    6  * * Redistributions of source code must retain the above copyright notice, this list of
    7  * conditions and the following disclaimer.
    8  * * Redistributions in binary form must reproduce the above copyright notice, this list of
    9  * conditions and the following disclaimer in the documentation and/or other materials
    10  * provided with the distribution.
    11  * * Neither the name of the NVIDIA CORPORATION nor the names of its contributors may be used
    12  * to endorse or promote products derived from this software without specific prior written
    13  * permission.
    14  *
    15  * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND ANY EXPRESS OR
    16  * IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND
    17  * FITNESS FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE
    18  * FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING,
    19  * BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS;
    20  * OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT,
    21  * STRICT LIABILITY, OR TOR (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
    22  * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
    23  *
    24  **************************************************************************************************/
    28 #pragma once
    29 
    30 namespace cutlass {
    31 
    33 
    35 struct MatrixLayout {
    36  enum Kind { kRowMajor, kColumnMajor };
    37 };
    38 
    40 
    42 struct GemmOperand {
    43  enum Kind { kA, kB, kC, kD };
    44 };
    45 
    47 
    48 } // namespace cutlass
    cutlass
    Definition: convert.h:33
    -
    cutlass::GemmOperand::kC
    Definition: matrix_traits.h:43
    -
    cutlass::MatrixLayout
    Describes layouts of matrices.
    Definition: matrix_traits.h:35
    -
    cutlass::MatrixLayout::kRowMajor
    Definition: matrix_traits.h:36
    -
    cutlass::GemmOperand::kB
    Definition: matrix_traits.h:43
    -
    cutlass::GemmOperand
    Gemm operand - D = A * B + C.
    Definition: matrix_traits.h:42
    -
    cutlass::MatrixLayout::kColumnMajor
    Definition: matrix_traits.h:36
    -
    cutlass::MatrixLayout::Kind
    Kind
    Definition: matrix_traits.h:36
    -
    cutlass::GemmOperand::Kind
    Kind
    Definition: matrix_traits.h:43
    -
    cutlass::GemmOperand::kA
    Definition: matrix_traits.h:43
    -
    cutlass::GemmOperand::kD
    Definition: matrix_traits.h:43
    +Go to the documentation of this file.
    1 /***************************************************************************************************
    2  * Copyright (c) 2017-2018, NVIDIA CORPORATION. All rights reserved.
    3  *
    4  * Redistribution and use in source and binary forms, with or without modification, are permitted
    5  * provided that the following conditions are met:
    6  * * Redistributions of source code must retain the above copyright notice, this list of
    7  * conditions and the following disclaimer.
    8  * * Redistributions in binary form must reproduce the above copyright notice, this list of
    9  * conditions and the following disclaimer in the documentation and/or other materials
    10  * provided with the distribution.
    11  * * Neither the name of the NVIDIA CORPORATION nor the names of its contributors may be used
    12  * to endorse or promote products derived from this software without specific prior written
    13  * permission.
    14  *
    15  * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND ANY EXPRESS OR
    16  * IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND
    17  * FITNESS FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE
    18  * FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING,
    19  * BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS;
    20  * OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT,
    21  * STRICT LIABILITY, OR TOR (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
    22  * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
    23  *
    24  **************************************************************************************************/
    28 #pragma once
    29 
    30 #include "cutlass/coord.h"
    31 
    32 namespace cutlass {
    33 
    35 
    38 struct MatrixCoord : public Coord<2, int> {
    39 
    41  typedef int Index;
    42 
    44  typedef Coord<2, Index> Base;
    45 
    47  static int const kRow = 0;
    48 
    50  static int const kColumn = 1;
    51 
    52  //
    53  // Methods
    54  //
    55 
    57  CUTLASS_HOST_DEVICE
    58  MatrixCoord() { }
    59 
    61  CUTLASS_HOST_DEVICE
    62  MatrixCoord(Coord<2, Index> const &coord): Base(coord) { }
    63 
    65  CUTLASS_HOST_DEVICE
    66  MatrixCoord(Index row, Index column): Base(make_Coord(row, column)) { }
    67 
    69  CUTLASS_HOST_DEVICE
    70  Index const & row() const { return this->at(kRow); }
    71 
    73  CUTLASS_HOST_DEVICE
    74  Index & row() { return this->at(kRow); }
    75 
    77  CUTLASS_HOST_DEVICE
    78  Index const & column() const { return this->at(kColumn); }
    79 
    81  CUTLASS_HOST_DEVICE
    82  Index & column() { return this->at(kColumn); }
    83 
    84  //
    85  // Coord operators
    86  //
    87 
    89  CUTLASS_HOST_DEVICE
    90  MatrixCoord operator+(Base const& b) const {
    91  return MatrixCoord(Base::operator+(b));
    92  }
    93 
    95  CUTLASS_HOST_DEVICE
    96  MatrixCoord operator-(Base const& b) const {
    97  return MatrixCoord(Base::operator-(b));
    98  }
    99 
    101  CUTLASS_HOST_DEVICE
    102  MatrixCoord operator*(Base const& b) const {
    103  return MatrixCoord(Base::operator*(b));
    104  }
    105 
    107  CUTLASS_HOST_DEVICE
    108  MatrixCoord operator/(Base const& b) const {
    109  return MatrixCoord(Base::operator/(b));
    110  }
    111 
    113  CUTLASS_HOST_DEVICE
    114  MatrixCoord& operator+=(Base const& b) {
    115  Base::operator+=(b);
    116  return *this;
    117  }
    118 
    120  CUTLASS_HOST_DEVICE
    121  MatrixCoord& operator-=(Base const& b) {
    122  Base::operator-=(b);
    123  return *this;
    124  }
    125 
    127  CUTLASS_HOST_DEVICE
    128  MatrixCoord& operator*=(Base const& b) {
    129  Base::operator*=(b);
    130  return *this;
    131  }
    132 
    134  CUTLASS_HOST_DEVICE
    135  MatrixCoord& operator/=(Base const& b) {
    136  Base::operator/=(b);
    137  return *this;
    138  }
    139 };
    140 
    142 
    144 //
    145 // The following define classes satisfying the TensorRefMapFunc concept. These must support the
    146 // following operations, where func is an instance of type TensorRefMapFunc.
    147 //
    148 // Coord<TensorRefMapFunc::kStorageRank> = func(Coord<kRank>);
    149 //
    150 // Though not required to be usable by TensorRef, each of the following also define a helper
    151 // function to map the "leading dimension" to an appropriate stride vector. Implementations
    152 // following this convention should also implement the following static method:
    153 //
    154 // Coord<TensorRefMapFunc::kStorageRank> stride = TensorRefMapFunc::stride(leading_dim);
    155 //
    156 struct MatrixLayout {
    157 
    159  enum Kind { kRowMajor, kColumnMajor };
    160 
    161  //
    162  // TensorRefMapFunc definitions for common layouts
    163  //
    164 
    166  struct RowMajor {
    167  static int const kStorageRank = 2;
    169  CUTLASS_HOST_DEVICE
    170  Coord<kStorageRank> operator()(MatrixCoord const &coord) const {
    171  return coord;
    172  }
    173  };
    174 
    176  struct ColumnMajor {
    177  static int const kStorageRank = 2;
    179  CUTLASS_HOST_DEVICE
    180  Coord<kStorageRank> operator()(MatrixCoord const &coord) const {
    181  return make_Coord(coord.column(), coord.row());
    182  }
    183  };
    184 
    187  template <int Interleave>
    188  struct RowMajorInterleaved {
    189 
    191  static int const kStorageRank = 3;
    192 
    194  static int const kInterleave = Interleave;
    195 
    197  CUTLASS_HOST_DEVICE
    198  Coord<kStorageRank> operator()(MatrixCoord const &coord) const {
    199  return make_Coord(
    200  coord.row() / kInterleave,
    201  coord.column(),
    202  coord.row() % kInterleave
    203  );
    204  }
    205 
    207  CUTLASS_HOST_DEVICE
    208  static Coord<kStorageRank> stride(int ldm) {
    209  return make_Coord(
    210  ldm * kInterleave,
    211  kInterleave,
    212  1
    213  );
    214  }
    215  };
    216 
    219  template <int Interleave>
    220  struct ColumnMajorInterleaved {
    221 
    223  static int const kStorageRank = 3;
    224 
    226  static int const kInterleave = Interleave;
    227 
    229  CUTLASS_HOST_DEVICE
    230  Coord<kStorageRank> operator()(MatrixCoord const &coord) const {
    231  return make_Coord(
    232  coord.column() / kInterleave,
    233  coord.row(),
    234  coord.column() % kInterleave
    235  );
    236  }
    237 
    239  CUTLASS_HOST_DEVICE
    240  static Coord<kStorageRank> stride(int ldm) {
    241  return make_Coord(
    242  ldm * kInterleave,
    243  kInterleave,
    244  1
    245  );
    246  }
    247  };
    248 
    251  struct ContiguousLayout {
    253  static int const kStorageRank = 3;
    254 
    256  static int const kRow = 0;
    257 
    259  static int const kColumn = 1;
    260 
    263  CUTLASS_HOST_DEVICE
    264  Coord<kStorageRank> operator()(MatrixCoord const &coord) const {
    265  return make_Coord(coord.row(), coord.column(), 0);
    266  }
    267 
    269  CUTLASS_HOST_DEVICE
    270  static Coord<kStorageRank> stride(MatrixLayout::Kind layout, int ldm) {
    271  if (layout == MatrixLayout::kRowMajor) {
    272  return make_Coord(ldm, 1, 1);
    273  }
    274  return make_Coord(1, ldm, 1);
    275  }
    276  };
    277 
    280  template <int BlockRows, int BlockColumns>
    281  struct ColumnMajorBlockLinear {
    282 
    284  static int const kStorageRank = 4;
    285 
    287  static int const kBlockRows = BlockRows;
    288 
    290  static int const kBlockColumns = BlockColumns;
    291 
    293  CUTLASS_HOST_DEVICE
    294  Coord<kStorageRank> operator()(MatrixCoord const &coord) const {
    295  return make_Coord(
    296  coord.column() / kBlockColumns,
    297  coord.row() / kBlockRows,
    298  coord.column() % kBlockColumns,
    299  coord.row() % kBlockRows
    300  );
    301  }
    302 
    304  CUTLASS_HOST_DEVICE
    305  static Coord<kStorageRank> stride(int ldm) {
    306  return make_Coord(
    307  ldm * kBlockRows * kBlockColumns,
    308  kBlockRows * kBlockColumns,
    309  kBlockRows,
    310  1
    311  );
    312  }
    313  };
    314 
    317  template <int BlockRows, int BlockColumns>
    318  struct RowMajorBlockLinear {
    319 
    321  static int const kStorageRank = 4;
    322 
    324  static int const kBlockRows = BlockRows;
    325 
    327  static int const kBlockColumns = BlockColumns;
    328 
    330  CUTLASS_HOST_DEVICE
    331  Coord<kStorageRank> operator()(MatrixCoord const &coord) const {
    332  return make_Coord(
    333  coord.row() / kBlockRows,
    334  coord.column() / kBlockColumns,
    335  coord.row() % kBlockRows,
    336  coord.column() % kBlockColumns
    337  );
    338  }
    339 
    341  CUTLASS_HOST_DEVICE
    342  static Coord<kStorageRank> stride(int ldm) {
    343  return make_Coord(
    344  ldm * kBlockRows * kBlockColumns,
    345  kBlockRows * kBlockColumns,
    346  kBlockColumns,
    347  1
    348  );
    349  }
    350  };
    351 };
    352 
    354 
    356 struct GemmOperand {
    357  enum Kind { kA, kB, kC, kD };
    358 };
    359 
    361 
    363 struct MatrixTransform {
    364  enum Kind {
    365  kNone,
    366  kConjugate,
    367  };
    368 };
    369 
    371 
    372 } // namespace cutlass
    cutlass::MatrixCoord::Index
    int Index
    Integer-valued index.
    Definition: matrix_traits.h:41
    +
    cutlass::MatrixLayout::ColumnMajor
    Mapping function for column-major matrices.
    Definition: matrix_traits.h:176
    +
    cutlass::MatrixLayout::RowMajorBlockLinear::kBlockColumns
    static int const kBlockColumns
    Interleaving size in columns dimension.
    Definition: matrix_traits.h:327
    +
    cutlass
    Definition: convert.h:33
    +
    cutlass::MatrixLayout::ColumnMajorInterleaved::operator()
    CUTLASS_HOST_DEVICE Coord< kStorageRank > operator()(MatrixCoord const &coord) const
    Maps (row, col) to (col, row, col)
    Definition: matrix_traits.h:230
    +
    cutlass::MatrixLayout::RowMajor::operator()
    CUTLASS_HOST_DEVICE Coord< kStorageRank > operator()(MatrixCoord const &coord) const
    Maps (i, j) to (i, j)
    Definition: matrix_traits.h:170
    +
    cutlass::MatrixTransform
    Transformation applied to matrix operands.
    Definition: matrix_traits.h:363
    +
    cutlass::MatrixLayout::RowMajorInterleaved
    Definition: matrix_traits.h:188
    +
    cutlass::MatrixLayout::ColumnMajorBlockLinear::kBlockColumns
    static int const kBlockColumns
    Interleaving size in columns dimension.
    Definition: matrix_traits.h:290
    +
    cutlass::MatrixTransform::kNone
    Definition: matrix_traits.h:365
    +
    cutlass::MatrixLayout::ColumnMajorBlockLinear
    Definition: matrix_traits.h:281
    +
    cutlass::MatrixLayout::ColumnMajorInterleaved
    Definition: matrix_traits.h:220
    +
    coord.h
    A Coord is a coordinate of arbitrary rank into a tensor or matrix.
    +
    cutlass::make_Coord
    CUTLASS_HOST_DEVICE Coord< 1 > make_Coord(int _0)
    Helper to make a 2-element coordinate.
    Definition: coord.h:318
    +
    cutlass::MatrixTransform::kConjugate
    no operation
    Definition: matrix_traits.h:366
    +
    cutlass::MatrixCoord::operator/=
    CUTLASS_HOST_DEVICE MatrixCoord & operator/=(Base const &b)
    In-place division.
    Definition: matrix_traits.h:135
    +
    cutlass::MatrixLayout::RowMajor::kStorageRank
    static int const kStorageRank
    Definition: matrix_traits.h:167
    +
    cutlass::MatrixLayout::ContiguousLayout
    Definition: matrix_traits.h:251
    +
    cutlass::MatrixLayout::ColumnMajor::operator()
    CUTLASS_HOST_DEVICE Coord< kStorageRank > operator()(MatrixCoord const &coord) const
    Maps (i, j) to (j, i)
    Definition: matrix_traits.h:180
    +
    cutlass::MatrixTransform::Kind
    Kind
    Definition: matrix_traits.h:364
    +
    cutlass::MatrixCoord::Base
    Coord< 2, Index > Base
    Base type is a Coord of rank=2.
    Definition: matrix_traits.h:44
    +
    cutlass::MatrixCoord::operator+
    CUTLASS_HOST_DEVICE MatrixCoord operator+(Base const &b) const
    Element-wise addition.
    Definition: matrix_traits.h:90
    +
    cutlass::Coord::operator*=
    CUTLASS_HOST_DEVICE Coord & operator*=(Coord const &b)
    In-place multiplication.
    Definition: coord.h:197
    +
    cutlass::GemmOperand::kC
    Definition: matrix_traits.h:357
    +
    cutlass::MatrixLayout::ContiguousLayout::kRow
    static int const kRow
    Dimension of rows.
    Definition: matrix_traits.h:256
    +
    cutlass::MatrixLayout::ColumnMajor::kStorageRank
    static int const kStorageRank
    Definition: matrix_traits.h:177
    +
    cutlass::MatrixLayout::ColumnMajorBlockLinear::kBlockRows
    static int const kBlockRows
    Interleaving size in rows dimension.
    Definition: matrix_traits.h:287
    +
    cutlass::MatrixLayout
    Defines data layouts of various matrix formats usable by TensorRef and other classes.
    Definition: matrix_traits.h:156
    +
    cutlass::MatrixLayout::RowMajorInterleaved::kInterleave
    static int const kInterleave
    Interleaving size.
    Definition: matrix_traits.h:194
    +
    cutlass::MatrixLayout::kRowMajor
    Definition: matrix_traits.h:159
    +
    cutlass::MatrixCoord::column
    CUTLASS_HOST_DEVICE Index const & column() const
    Returns the column of the coordinate.
    Definition: matrix_traits.h:78
    +
    cutlass::MatrixCoord::MatrixCoord
    CUTLASS_HOST_DEVICE MatrixCoord(Index row, Index column)
    Helper to construct from a row and column.
    Definition: matrix_traits.h:66
    +
    cutlass::MatrixLayout::RowMajorInterleaved::stride
    static CUTLASS_HOST_DEVICE Coord< kStorageRank > stride(int ldm)
    Helper to compute stride vector from leading dimension.
    Definition: matrix_traits.h:208
    +
    cutlass::MatrixLayout::ContiguousLayout::kColumn
    static int const kColumn
    Dimension of columns.
    Definition: matrix_traits.h:259
    +
    cutlass::MatrixLayout::RowMajorInterleaved::kStorageRank
    static int const kStorageRank
    Rank of storage n-D array.
    Definition: matrix_traits.h:191
    +
    cutlass::Coord::operator-=
    CUTLASS_HOST_DEVICE Coord & operator-=(Coord const &b)
    In-place subtraction.
    Definition: coord.h:188
    +
    cutlass::MatrixLayout::ContiguousLayout::kStorageRank
    static int const kStorageRank
    Arbitrary storage rank.
    Definition: matrix_traits.h:253
    +
    cutlass::GemmOperand::kB
    Definition: matrix_traits.h:357
    +
    cutlass::Coord::operator+=
    CUTLASS_HOST_DEVICE Coord & operator+=(Coord const &b)
    In-place addition.
    Definition: coord.h:179
    +
    cutlass::MatrixLayout::RowMajorInterleaved::operator()
    CUTLASS_HOST_DEVICE Coord< kStorageRank > operator()(MatrixCoord const &coord) const
    Maps (row, col) to (row, col, row)
    Definition: matrix_traits.h:198
    +
    CUTLASS_HOST_DEVICE
    #define CUTLASS_HOST_DEVICE
    Definition: cutlass.h:46
    +
    cutlass::MatrixLayout::RowMajorBlockLinear::kBlockRows
    static int const kBlockRows
    Interleaving size in rows dimension.
    Definition: matrix_traits.h:324
    +
    cutlass::MatrixCoord::row
    CUTLASS_HOST_DEVICE Index const & row() const
    Returns the row of the coordinate.
    Definition: matrix_traits.h:70
    +
    cutlass::Coord< 2, int >::at
    CUTLASS_HOST_DEVICE Index & at()
    Gets the index of a given Coord element.
    Definition: coord.h:240
    +
    cutlass::Coord::operator/=
    CUTLASS_HOST_DEVICE Coord & operator/=(Coord const &b)
    In-place division.
    Definition: coord.h:206
    +
    cutlass::MatrixCoord::operator-
    CUTLASS_HOST_DEVICE MatrixCoord operator-(Base const &b) const
    Element-wise subtraction.
    Definition: matrix_traits.h:96
    +
    cutlass::MatrixCoord::MatrixCoord
    CUTLASS_HOST_DEVICE MatrixCoord(Coord< 2, Index > const &coord)
    Constructs from Coord<2>
    Definition: matrix_traits.h:62
    +
    cutlass::MatrixLayout::RowMajorBlockLinear::kStorageRank
    static int const kStorageRank
    Rank of storage n-D array.
    Definition: matrix_traits.h:321
    +
    cutlass::Coord
    Statically-sized array specifying Coords within a tensor.
    Definition: coord.h:49
    +
    cutlass::GemmOperand
    Gemm operand - D = A * B + C.
    Definition: matrix_traits.h:356
    +
    cutlass::MatrixLayout::RowMajorBlockLinear::stride
    static CUTLASS_HOST_DEVICE Coord< kStorageRank > stride(int ldm)
    Helper to compute stride vector from leading dimension.
    Definition: matrix_traits.h:342
    +
    cutlass::MatrixCoord::kRow
    static int const kRow
    Rows dimension.
    Definition: matrix_traits.h:47
    +
    cutlass::MatrixCoord::operator-=
    CUTLASS_HOST_DEVICE MatrixCoord & operator-=(Base const &b)
    In-place subtraction.
    Definition: matrix_traits.h:121
    +
    cutlass::MatrixCoord::operator*
    CUTLASS_HOST_DEVICE MatrixCoord operator*(Base const &b) const
    Element-wise multiplication.
    Definition: matrix_traits.h:102
    +
    cutlass::MatrixLayout::kColumnMajor
    Definition: matrix_traits.h:159
    +
    cutlass::MatrixLayout::ColumnMajorInterleaved::stride
    static CUTLASS_HOST_DEVICE Coord< kStorageRank > stride(int ldm)
    Helper to compute stride vector from leading dimension.
    Definition: matrix_traits.h:240
    +
    cutlass::MatrixLayout::ContiguousLayout::operator()
    CUTLASS_HOST_DEVICE Coord< kStorageRank > operator()(MatrixCoord const &coord) const
    Definition: matrix_traits.h:264
    +
    cutlass::MatrixLayout::Kind
    Kind
    Enumeration defining fundamental contiguous layouts.
    Definition: matrix_traits.h:159
    +
    cutlass::MatrixCoord::row
    CUTLASS_HOST_DEVICE Index & row()
    Returns the row of the coordinate.
    Definition: matrix_traits.h:74
    +
    cutlass::MatrixLayout::ColumnMajorBlockLinear::kStorageRank
    static int const kStorageRank
    Rank of storage n-D array.
    Definition: matrix_traits.h:284
    +
    cutlass::MatrixLayout::ColumnMajorInterleaved::kInterleave
    static int const kInterleave
    Interleaving size.
    Definition: matrix_traits.h:226
    +
    cutlass::MatrixLayout::RowMajorBlockLinear::operator()
    CUTLASS_HOST_DEVICE Coord< kStorageRank > operator()(MatrixCoord const &coord) const
    Maps (row, col) to (row, col, row, col)
    Definition: matrix_traits.h:331
    +
    cutlass::GemmOperand::Kind
    Kind
    Definition: matrix_traits.h:357
    +
    cutlass::GemmOperand::kA
    Definition: matrix_traits.h:357
    +
    cutlass::MatrixCoord::column
    CUTLASS_HOST_DEVICE Index & column()
    Returns the column of the coordinate.
    Definition: matrix_traits.h:82
    +
    cutlass::MatrixCoord::operator*=
    CUTLASS_HOST_DEVICE MatrixCoord & operator*=(Base const &b)
    In-place multiplication.
    Definition: matrix_traits.h:128
    +
    cutlass::MatrixLayout::ColumnMajorInterleaved::kStorageRank
    static int const kStorageRank
    Rank of storage n-D array.
    Definition: matrix_traits.h:223
    +
    cutlass::MatrixLayout::RowMajorBlockLinear
    Definition: matrix_traits.h:318
    +
    cutlass::MatrixCoord::operator+=
    CUTLASS_HOST_DEVICE MatrixCoord & operator+=(Base const &b)
    In-place addition.
    Definition: matrix_traits.h:114
    +
    cutlass::MatrixLayout::ColumnMajorBlockLinear::stride
    static CUTLASS_HOST_DEVICE Coord< kStorageRank > stride(int ldm)
    Helper to compute stride vector from leading dimension.
    Definition: matrix_traits.h:305
    +
    cutlass::MatrixLayout::ColumnMajorBlockLinear::operator()
    CUTLASS_HOST_DEVICE Coord< kStorageRank > operator()(MatrixCoord const &coord) const
    Maps (row, col) to (col, row, col, row)
    Definition: matrix_traits.h:294
    +
    cutlass::MatrixLayout::ContiguousLayout::stride
    static CUTLASS_HOST_DEVICE Coord< kStorageRank > stride(MatrixLayout::Kind layout, int ldm)
    Helper to construct a stride vector based on contiguous matrix layout and leading dimension...
    Definition: matrix_traits.h:270
    +
    cutlass::MatrixCoord
    Definition: matrix_traits.h:38
    +
    cutlass::MatrixCoord::operator/
    CUTLASS_HOST_DEVICE MatrixCoord operator/(Base const &b) const
    Element-wise division.
    Definition: matrix_traits.h:108
    +
    cutlass::MatrixCoord::kColumn
    static int const kColumn
    Columns dimension.
    Definition: matrix_traits.h:50
    +
    cutlass::MatrixCoord::MatrixCoord
    CUTLASS_HOST_DEVICE MatrixCoord()
    Default ctor.
    Definition: matrix_traits.h:58
    +
    cutlass::GemmOperand::kD
    Definition: matrix_traits.h:357
    +
    cutlass::MatrixLayout::RowMajor
    Mapping function for row-major matrices.
    Definition: matrix_traits.h:166
    -Generated on Thu May 3 2018 16:36:12 for Cutlass by   +Generated on Wed Sep 19 2018 13:58:50 for Cutlass by   doxygen 1.8.14
    diff --git a/docs/menudata.js b/docs/menudata.js index 725988aa..dde1bbfe 100644 --- a/docs/menudata.js +++ b/docs/menudata.js @@ -29,24 +29,33 @@ var menudata={children:[ {text:"Namespace Members",url:"namespacemembers.html",children:[ {text:"All",url:"namespacemembers.html",children:[ {text:"_",url:"namespacemembers.html#index__"}, +{text:"a",url:"namespacemembers.html#index_a"}, {text:"c",url:"namespacemembers.html#index_c"}, +{text:"e",url:"namespacemembers.html#index_e"}, {text:"f",url:"namespacemembers.html#index_f"}, {text:"g",url:"namespacemembers.html#index_g"}, {text:"i",url:"namespacemembers.html#index_i"}, {text:"l",url:"namespacemembers.html#index_l"}, {text:"m",url:"namespacemembers.html#index_m"}, +{text:"n",url:"namespacemembers.html#index_n"}, {text:"o",url:"namespacemembers.html#index_o"}, +{text:"p",url:"namespacemembers.html#index_p"}, {text:"r",url:"namespacemembers.html#index_r"}, {text:"s",url:"namespacemembers.html#index_s"}, {text:"t",url:"namespacemembers.html#index_t"}]}, {text:"Functions",url:"namespacemembers_func.html",children:[ {text:"_",url:"namespacemembers_func.html#index__"}, +{text:"a",url:"namespacemembers_func.html#index_a"}, {text:"c",url:"namespacemembers_func.html#index_c"}, +{text:"e",url:"namespacemembers_func.html#index_e"}, +{text:"f",url:"namespacemembers_func.html#index_f"}, {text:"g",url:"namespacemembers_func.html#index_g"}, {text:"i",url:"namespacemembers_func.html#index_i"}, {text:"l",url:"namespacemembers_func.html#index_l"}, {text:"m",url:"namespacemembers_func.html#index_m"}, +{text:"n",url:"namespacemembers_func.html#index_n"}, {text:"o",url:"namespacemembers_func.html#index_o"}, +{text:"p",url:"namespacemembers_func.html#index_p"}, {text:"r",url:"namespacemembers_func.html#index_r"}, {text:"s",url:"namespacemembers_func.html#index_s"}]}, {text:"Typedefs",url:"namespacemembers_type.html"}]}]}, @@ -78,6 +87,7 @@ var menudata={children:[ {text:"v",url:"functions_v.html#index_v"}, {text:"w",url:"functions_w.html#index_w"}, {text:"y",url:"functions_y.html#index_y"}, +{text:"z",url:"functions_z.html#index_z"}, {text:"~",url:"functions_0x7e.html#index_0x7e"}]}, {text:"Functions",url:"functions_func.html",children:[ {text:"a",url:"functions_func.html#index_a"}, @@ -89,8 +99,10 @@ var menudata={children:[ {text:"g",url:"functions_func_g.html#index_g"}, {text:"h",url:"functions_func_h.html#index_h"}, {text:"i",url:"functions_func_i.html#index_i"}, +{text:"k",url:"functions_func_k.html#index_k"}, {text:"l",url:"functions_func_l.html#index_l"}, {text:"m",url:"functions_func_m.html#index_m"}, +{text:"n",url:"functions_func_n.html#index_n"}, {text:"o",url:"functions_func_o.html#index_o"}, {text:"p",url:"functions_func_p.html#index_p"}, {text:"r",url:"functions_func_r.html#index_r"}, @@ -99,6 +111,7 @@ var menudata={children:[ {text:"u",url:"functions_func_u.html#index_u"}, {text:"v",url:"functions_func_v.html#index_v"}, {text:"w",url:"functions_func_w.html#index_w"}, +{text:"z",url:"functions_func_z.html#index_z"}, {text:"~",url:"functions_func_0x7e.html#index_0x7e"}]}, {text:"Variables",url:"functions_vars.html",children:[ {text:"a",url:"functions_vars.html#index_a"}, @@ -113,6 +126,7 @@ var menudata={children:[ {text:"l",url:"functions_vars_l.html#index_l"}, {text:"m",url:"functions_vars_m.html#index_m"}, {text:"n",url:"functions_vars_n.html#index_n"}, +{text:"o",url:"functions_vars_o.html#index_o"}, {text:"p",url:"functions_vars_p.html#index_p"}, {text:"r",url:"functions_vars_r.html#index_r"}, {text:"s",url:"functions_vars_s.html#index_s"}, @@ -127,6 +141,7 @@ var menudata={children:[ {text:"f",url:"functions_type_f.html#index_f"}, {text:"g",url:"functions_type_g.html#index_g"}, {text:"i",url:"functions_type_i.html#index_i"}, +{text:"k",url:"functions_type_k.html#index_k"}, {text:"l",url:"functions_type_l.html#index_l"}, {text:"m",url:"functions_type_m.html#index_m"}, {text:"n",url:"functions_type_n.html#index_n"}, @@ -140,8 +155,10 @@ var menudata={children:[ {text:"Enumerations",url:"functions_enum.html"}, {text:"Enumerator",url:"functions_eval.html",children:[ {text:"a",url:"functions_eval.html#index_a"}, +{text:"b",url:"functions_eval.html#index_b"}, {text:"k",url:"functions_eval.html#index_k"}, {text:"m",url:"functions_eval.html#index_m"}, +{text:"o",url:"functions_eval.html#index_o"}, {text:"v",url:"functions_eval.html#index_v"}]}]}]}, {text:"Files",url:"files.html",children:[ {text:"File List",url:"files.html"}, diff --git a/docs/modules.html b/docs/modules.html index c42247bd..8fc90844 100644 --- a/docs/modules.html +++ b/docs/modules.html @@ -76,19 +76,20 @@ $(function() {

    Namespaces

    - - - - - - - + + + + + + + +
     Fragment Concept
     Fragment Iterator Concept
     Predicate Vector Concept
     Predicate Iterator Concept
     Predicate Tile Adapter Concept
     Layout Concept
     Tile Traits Concept
     Tile Load Iterator Concept
     Tile Store Iterator Concept
     Identity Block Swizzle
     Predicate Vector Concept
     Predicate Iterator Concept
     Predicate Tile Adapter Concept
     Layout Concept
     Tile Traits Concept
     Tile Load Iterator Concept
     Tile Store Iterator Concept
    -Generated on Thu May 3 2018 16:36:15 for Cutlass by   +Generated on Wed Sep 19 2018 13:58:53 for Cutlass by   doxygen 1.8.14
    diff --git a/docs/namespacecutlass.html b/docs/namespacecutlass.html index 989135cb..4fb1ce9a 100644 --- a/docs/namespacecutlass.html +++ b/docs/namespacecutlass.html @@ -79,6 +79,8 @@ $(function() { + + @@ -88,24 +90,14 @@ Namespaces Classes + + - - - - - - - - - - - - @@ -129,6 +121,8 @@ Classes + + @@ -143,44 +137,50 @@ Classes + + + - - - - - - - - - - - - + + + + - - + + - + - + + - + - + + + + + + + + + + + @@ -190,22 +190,43 @@ Classes + + - + + + + + + + + + + + + + + + + + + + + @@ -216,6 +237,8 @@ Classes + + @@ -240,19 +263,34 @@ Classes - + - + - + - + + + + + - + + + + + + + - + + + + + + @@ -263,9 +301,15 @@ Classes + + + + + + @@ -282,13 +326,30 @@ Classes + + + + + + + + + + + + + - + + + + + @@ -299,6 +360,20 @@ Classes + + + + + + + + + + + + + +

    Namespaces

     detail
     
     gemm
     
     platform
    struct  AlignedStruct
     
    struct  bin1_t
     
    struct  ComputeOffsetFromShape
     Compute the offset for the given coordinates in a cube. More...
     
    struct  ComputeOffsetFromShape< Shape< 1, kSh_, kSw_, 1 > >
     Compute the offset for the given coordinates in a cube with one channel and a depth of 1. More...
     
    struct  ComputeOffsetFromShape< Shape< 1, kSh_, kSw_, kSc_ > >
     Compute the offset for the given coordinates in a cube with a depth of 1. More...
     
    struct  ComputeOffsetFromStrides
     Compute the offset for the given coordinates in a cube. More...
     
    struct  ComputeOffsetFromStrides< Shape< 1, S_h_, S_w_, 1 > >
     Compute the offset for the given coordinates in a cube with one channel and a depth of 1. More...
     
    struct  ComputeOffsetFromStrides< Shape< 1, S_h_, S_w_, S_c_ > >
     Compute the offset for the given coordinates in a cube with a depth of 1. More...
     
    struct  ComputeThreadOffsetFromStrides
     Decompose threadId.x into coordinate of a cube whose dimensions are specified by Threads_. Afterwards compute the offset of those coordinates using Strides_. More...
     
     
    struct  divide_assert
     
    struct  DumpType
     
    struct  Extent
     Returns the extent of a scalar or vector. More...
     
     
    struct  FragmentConstIterator
     
    struct  FragmentElementType
     Specifies whether iterator storage fragment consists of Scalar values or WMMA matrix. More...
     
    struct  FragmentIterator
     A template defining Fragment Iterator Concept. More...
     
    struct  FragmentLoad
     
    struct  FragmentLoad< IteratorFragment::kScalar, kAccessSize, Scalar_, Memory_, FragmentElement_, kStride >
     
    struct  FragmentLoad< IteratorFragment::kWmmaMatrix, kAccessSize, Scalar_, Memory_, FragmentElement_, kStride >
     
    struct  FragmentStore
     
    struct  FragmentStore< IteratorFragment::kScalar, kAccessSize, Scalar_, Memory_, FragmentElement_, kStride >
     
    struct  FragmentStore< IteratorFragment::kWmmaMatrix, kAccessSize, Scalar_, Memory_, FragmentElement_, kStride >
     
    struct  GemmOperand
     Gemm operand - D = A * B + C. More...
     
    struct  Identity
     Describes identity elements. More...
     
    struct  IdentityTensorMapFunc
     
    struct  int4_t
     
    struct  is_pow2
     
    struct  IteratorAdvance
     Specifies dimension in which post-increment accesses advance. More...
     
    struct  IteratorFragment
     Specifies whether iterator storage fragment consists of Scalar values or WMMA matrix. More...
    struct  KernelLaunchConfiguration
     Structure containing the basic launch configuration of a CUDA kernel. More...
     
    struct  Load
     
    struct  Load< double, 2, Memory_, true, 16 >
    struct  Load< double, 2, Memory_, FragmentElementType::kScalar, double, kStride, 16 >
     
    struct  Load< Scalar_, Lanes_, Memory_, true, 16 >
    struct  Load< Scalar_, kAccessSize, Memory_, FragmentElementType::kScalar, Scalar_, 1, 2 >
     Partial specialization for 16b loads. More...
     
    struct  Load< Scalar_, Lanes_, Memory_, true, 4 >
    struct  Load< Scalar_, kAccessSize, Memory_, FragmentElementType::kScalar, Scalar_, kStride, 16 >
     
    struct  Load< Scalar_, Lanes_, Memory_, true, 8 >
    struct  Load< Scalar_, kAccessSize, Memory_, FragmentElementType::kScalar, Scalar_, kStride, 4 >
     
    struct  Load< Scalar_, kAccessSize, Memory_, FragmentElementType::kScalar, Scalar_, kStride, 8 >
     
    struct  Load< Scalar_, kAccessSize, Memory_, FragmentElementType::kWmmaMatrix, FragmentElement_, kStride, size >
     
    struct  Load< Vector< bin1_t, 32 >, kAccessSize, Memory_, FragmentElementType::kWmmaMatrix, FragmentElement_, kStride, size >
     
    struct  Load< Vector< int4_t, 8 >, kAccessSize, Memory_, FragmentElementType::kWmmaMatrix, FragmentElement_, kStride, size >
     
    struct  Load< Vector< uint4_t, 8 >, kAccessSize, Memory_, FragmentElementType::kWmmaMatrix, FragmentElement_, kStride, size >
     
    struct  log2_down
     
     
    struct  log2_up< N, 1, Count >
     
    struct  MatrixCoord
     
    struct  MatrixLayout
     Describes layouts of matrices. More...
     Defines data layouts of various matrix formats usable by TensorRef and other classes. More...
     
    struct  MatrixTransform
     Transformation applied to matrix operands. More...
     
    struct  Max
     
    struct  MemorySpace
     Enum to specify which memory space data resides in. More...
     
    struct  Min
     
    struct  PredicatedTileLoadStream
     Generic stream for loading and transforming fragments. More...
     
    struct  PredicatedTileStoreStream
     Generic stream for transforming and storing fragments. More...
     
    struct  PredicateTileAdapter
     Adapter to enable random access to predicates via logical coordinate within a tile. More...
     
    struct  PredicateVector
     Statically sized array of bits implementing. More...
     
    struct  RegularTilePredicateFunctor
     Functor computing a predicate given the logical position of an access. More...
     
    struct  ReshapeTile
     
    struct  ReshapeTile< Tile_, kAccessSize_, true >
     
    struct  ScalarIO
     Helper to enable formatted printing of CUTLASS scalar types to an ostream. More...
     
    struct  Shape
     A Shape implementing Layout Concept describing the dimensions of a cube. More...
     
     
    struct  ShapeDiv
     
    struct  ShapeDivCeiling
     
    struct  ShapeMax
     
    struct  ShapeMin
     
    struct  Store
     
    struct  Store< double, 2, Memory_, true, 16 >
    struct  Store< double, 2, Memory_, FragmentElementType::kScalar, double, kStride, 16 >
     
    struct  Store< Scalar_, Lanes_, Memory_, true, 16 >
    struct  Store< Scalar_, kAccessSize, Memory_, FragmentElementType::kScalar, Scalar_, 1, 2 >
     
    struct  Store< Scalar_, Lanes_, Memory_, true, 4 >
    struct  Store< Scalar_, kAccessSize, Memory_, FragmentElementType::kScalar, Scalar_, kStride, 16 >
     
    struct  Store< Scalar_, Lanes_, Memory_, true, 8 >
    struct  Store< Scalar_, kAccessSize, Memory_, FragmentElementType::kScalar, Scalar_, kStride, 4 >
     
    struct  Store< Scalar_, kAccessSize, Memory_, FragmentElementType::kScalar, Scalar_, kStride, 8 >
     
    struct  Store< Scalar_, kAccessSize, Memory_, FragmentElementType::kWmmaMatrix, FragmentElement_, kStride, size >
     
    class  TensorRef
     Structure modeling a pointer and stride into a tensor. More...
     
    class  TensorRef< Storage_, Rank_, MapFunc_, 1, Index_, LongIndex_ >
     Specialization for rank=1 case with no internal StrideVector. More...
     
    struct  TensorRefArray
     
    struct  TensorRefBatchStrided
     
    class  TensorView
     Host-side reference implementation of tensor operations. More...
     Defines a view into a logical tensor. More...
     
    struct  TileAllocation
     Class for storing a tile in memory and accessing it through a tensor ref. More...
     
    struct  TileCoord
     
    struct  TiledThreadOffset
     Basic thread offset function computed from a thread shape. More...
    struct  TileLoadIterator
     An iterator implementing Tile Load Iterator Concept for loading a tile from memory. More...
     
    struct  TileLoadStream
     Generic stream for loading and transforming fragments. More...
     
    struct  TileStoreIterator
     An iterator implementing Tile Store Iterator Concept for storing a tile to memory. More...
     
    struct  TileStoreStream
     Generic stream for transforming and storing fragments. More...
     
    struct  TileTraits
     A template defining Tile Traits Concept. More...
     
    struct  TrivialPredicateTileAdapter
     Always returns true predicate. More...
     
    struct  uint4_t
     
    union  Vector
     
    union  Vector< bin1_t, kLanes_ >
     Vector definition for 1-bit binary datatype. More...
     
    union  Vector< half, 1 >
     
    union  Vector< half, kLanes_ >
     
    union  Vector< int4_t, kLanes_ >
     Vector definition for 4-bit signed integer datatype. More...
     
    union  Vector< uint4_t, kLanes_ >
     Vector definition for 4-bit unsigned integer datatype. More...
     
    struct  Vectorize
     
    struct  Vectorize< Element_, 1 >
    struct  Vectorize< Vector< bin1_t, 32 >, kLanes_ >
     
    struct  Vectorize< Vector< int4_t, 8 >, kLanes_ >
     
    struct  Vectorize< Vector< uint4_t, 8 >, kLanes_ >
     
    struct  VectorTraits
     Traits describing properties of vectors and scalar-as-vectors. More...
    struct  VectorTraits< Vector< T, Lanes > const >
     Partial specialization for actual cutlass::Vector. More...
     
    struct  ZipConvert
     Zips two convert operations. More...
     
    struct  ZipFragment
     A template defining Fragment Concept. More...
     
    struct  ZipTensorRef
     
    struct  ZipTileAllocation
     Manages a pair of tile allocations as if they are one allocation. More...
     
    class  ZipTileIterator
     Constructs an iterator from a pair of iterators. More...
     
    @@ -314,85 +389,42 @@ Functions - - - - - - - - - - - - + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - @@ -402,6 +434,12 @@ Functions + + + + + + @@ -426,12 +464,24 @@ Functions - - - - - - + + + + + + + + + + + + + + + + + +

    Functions

    CUTLASS_HOST_DEVICE Coord< 4 > make_Coord (int _0, int _1, int _2, int _3)
     Helper to make a 4-element coordinate. More...
     
    CUTLASS_HOST_DEVICE Coord< 2 > get_Coord_hw (Coord< 3 > const &coord)
     Getter. More...
     
    CUTLASS_HOST_DEVICE Coord< 2 > get_Coord_hw (Coord< 4 > const &coord)
     Getter. More...
     
    CUTLASS_HOST_DEVICE Coord< 3 > get_Coord_hwc (Coord< 4 > const &coord)
     Getter. More...
     
    CUTLASS_HOST_DEVICE Coord< 3 > get_Coord_dhw (Coord< 4 > const &coord)
     Getter. More...
     
    template<typename Shape_ >
    CUTLASS_HOST_DEVICE Coord< 3 > make_Coord_from_shape ()
     
    template<int Rank>
    std::ostream & operator<< (std::ostream &out, Coord< Rank > const &coord)
     
    template<typename T >
    std::ostream & operator<< (std::ostream &out, ScalarIO< T > const &scalar)
     Default printing to ostream. More...
     
    template<>
    std::ostream & operator<< (std::ostream &out, ScalarIO< int8_t > const &scalar)
     Printing to ostream of int8_t as integer rather than character. More...
     
    template<>
    std::ostream & operator<< (std::ostream &out, ScalarIO< uint8_t > const &scalar)
     Printing to ostream of uint8_t as integer rather than character. More...
     
    template<>
    std::ostream & operator<< (std::ostream &out, ScalarIO< cutlass::Vector< cutlass::bin1_t, 32 > > const &scalar)
     Printing to ostream of vector of 1b elements. More...
     
    template<>
    std::ostream & operator<< (std::ostream &out, ScalarIO< cutlass::Vector< cutlass::int4_t, 8 > > const &scalar)
     Printing to ostream of vector of 4b signed integer elements. More...
     
    template<>
    std::ostream & operator<< (std::ostream &out, ScalarIO< cutlass::Vector< cutlass::uint4_t, 8 > > const &scalar)
     Printing to ostream of vector of 4b unsigned integer elements. More...
     
    template<typename InputIterator , typename Fragment >
    CUTLASS_HOST_DEVICE void iterator_load (InputIterator &iterator, Fragment &fragment)
     Loads a fragment from an input iterator. More...
     
    template<typename InputIterator , typename Fragment >
    CUTLASS_DEVICE void shared_iterator_load (InputIterator &iterator, Fragment &fragment)
     Loads a fragment from a shared memory input iterator. More...
     
    template<typename InputIterator , typename Fragment >
    CUTLASS_DEVICE void shared_iterator_load (InputIterator &iterator, Fragment &fragment, int d)
     Loads a fragment from a shared memory input iterator. More...
     
    template<typename InputIterator , typename Fragment , typename ConstPredicateAdapter >
    CUTLASS_HOST_DEVICE void iterator_load_post_increment (InputIterator &iterator, Fragment &fragment, typename InputIterator::Index offset, ConstPredicateAdapter predicate_adapter)
     Loads a fragment from an input iterator, masked by a predicate iterator. More...
     
    template<typename InputIterator , typename Fragment >
    CUTLASS_HOST_DEVICE void iterator_load_post_increment (InputIterator &iterator, Fragment &fragment, typename InputIterator::Index offset=0)
     Loads a fragment from an input iterator. More...
     
    template<typename InputIterator , typename Fragment , typename ConstPredicateAdapter >
    CUTLASS_HOST_DEVICE void iterator_load_post_increment (InputIterator &iterator, Fragment &fragment, ConstPredicateAdapter pred_it)
     Loads a fragment from an input iterator. More...
     
    template<typename InputIterator , typename Fragment , typename ConstPredicateAdapter >
    CUTLASS_HOST_DEVICE void iterator_load (InputIterator const &_iterator, Fragment &fragment, typename InputIterator::Index offset, ConstPredicateAdapter predicate_adapter)
     
    template<typename InputIterator , typename Fragment >
    CUTLASS_HOST_DEVICE void iterator_load (InputIterator const &iterator, Fragment &fragment, typename InputIterator::Index offset=0)
     Loads a fragment from an input iterator. More...
     
    template<typename InputIterator , typename Fragment , typename ConstPredicateAdapter >
    CUTLASS_HOST_DEVICE void iterator_load (InputIterator const &iterator, Fragment &fragment, ConstPredicateAdapter pred_it)
     Loads a fragment from an input iterator. More...
     
    template<typename OutputIterator , typename Fragment >
    CUTLASS_HOST_DEVICE void iterator_store (OutputIterator &iterator, Fragment &fragment)
     Stores a fragment to an output iterator. More...
     
    template<typename OutputIterator , typename Fragment >
    CUTLASS_DEVICE void shared_iterator_store (OutputIterator &iterator, Fragment const &fragment)
     Stores a fragment to a shared memory output iterator. More...
     
    template<typename OutputIterator , typename Fragment , typename ConstPredicateAdapter >
    CUTLASS_HOST_DEVICE void iterator_store_post_increment (OutputIterator &iterator, Fragment const &fragment, typename OutputIterator::Index offset, ConstPredicateAdapter predicate_adapter)
     Stores a fragment to an output iterator, masked by a predicate iterator. More...
     
    template<typename OutputIterator , typename Fragment >
    CUTLASS_HOST_DEVICE void iterator_store_post_increment (OutputIterator &iterator, Fragment const &fragment, typename OutputIterator::Index offset=0)
     Stores a fragment to an output iterator. More...
     
    template<typename OutputIterator , typename Fragment , typename ConstPredicateAdapter >
    CUTLASS_HOST_DEVICE void iterator_store_post_increment (OutputIterator &iterator, Fragment const &fragment, ConstPredicateAdapter pred_it)
     Stores a fragment to an output iterator. More...
     
    template<typename OutputIterator , typename Fragment , typename ConstPredicateAdapter >
    CUTLASS_HOST_DEVICE void iterator_store (OutputIterator const &_iterator, Fragment const &fragment, typename OutputIterator::Index offset, ConstPredicateAdapter predicate_adapter)
     Stores a fragment to an output iterator, masked by a predicate iterator. More...
     
    template<typename OutputIterator , typename Fragment >
    CUTLASS_HOST_DEVICE void iterator_store (OutputIterator const &iterator, Fragment const &fragment, typename OutputIterator::Index offset=0)
     Stores a fragment to an output iterator. More...
     
    template<typename OutputIterator , typename Fragment , typename ConstPredicateAdapter >
    CUTLASS_HOST_DEVICE void iterator_store (OutputIterator const &iterator, Fragment const &fragment, ConstPredicateAdapter pred_it)
     Stores a fragment to an output iterator. More...
     
    template<typename dividend_t , typename divisor_t >
    CUTLASS_HOST_DEVICE dividend_t round_nearest (dividend_t dividend, divisor_t divisor)
     
    template<typename value_t >
    CUTLASS_HOST_DEVICE value_t lcm (value_t a, value_t b)
     
    template<typename value_t >
    CUTLASS_HOST_DEVICE value_t clz (value_t x)
     
    template<typename value_t >
    CUTLASS_HOST_DEVICE value_t find_log2 (value_t x)
     
    __host__ CUTLASS_DEVICE cudaError_t cuda_perror_impl (cudaError_t error, const char *filename, int line)
     The corresponding error message is printed to stderr (or stdout in device code) along with the supplied source context. More...
     
    template<>
    struct __align__ (64) AlignedStruct< 64 >
     
    template<typename Scalar_ >
    CUTLASS_DEVICE void make_zero (Scalar_ &x)
     
    template<typename Scalar_ , int kLanes_>
    CUTLASS_DEVICE void make_zero (Vector< Scalar_, kLanes_ > &vec)
     
    template<typename Scalar_ >
    CUTLASS_HOST_DEVICE void make_zero (Scalar_ &x)
     
    template<typename Scalar_ , int kLanes_>
    CUTLASS_HOST_DEVICE void make_zero (Vector< Scalar_, kLanes_ > &vec)
     
    template<typename First , typename Second >
    CUTLASS_HOST_DEVICE ZipFragment< First, Second > make_ZipFragment (First const &first, Second const &second)
     Helper to construct a ZipFragment object. More...
     
    template<typename First , typename Second >
    CUTLASS_HOST_DEVICE ZipConvert< First, Second > make_ZipConvert (First const &first, Second const &second)
     Helper to construct a ZipConvert object. More...
     
    template<typename First , typename Second >
    CUTLASS_HOST_DEVICE ZipTensorRef< First, Second > make_ZipTensorRef (First const &first, Second const &second)
     Constructs a ZipTensorRef. More...
     

    Function Documentation

    @@ -452,10 +502,30 @@ template<>
    +
    + + +

    ◆ __align__() [2/7]

    + +
    +
    +
    +template<>
    + + + + + + + + +
    struct cutlass::__align__ (2 )
    +
    +
    -

    ◆ __align__() [2/7]

    +

    ◆ __align__() [3/7]

    @@ -475,7 +545,7 @@ template<>
    -

    ◆ __align__() [3/7]

    +

    ◆ __align__() [4/7]

    @@ -495,7 +565,7 @@ template<>
    -

    ◆ __align__() [4/7]

    +

    ◆ __align__() [5/7]

    @@ -515,7 +585,7 @@ template<>
    -

    ◆ __align__() [5/7]

    +

    ◆ __align__() [6/7]

    @@ -535,7 +605,7 @@ template<>
    -

    ◆ __align__() [6/7]

    +

    ◆ __align__() [7/7]

    @@ -554,23 +624,24 @@ template<>
    - -

    ◆ __align__() [7/7]

    + +

    ◆ clz()

    -template<>
    +template<typename value_t >
    - + - - + +
    struct cutlass::__align__ CUTLASS_HOST_DEVICE value_t cutlass::clz (2 )value_t x)
    +

    log2 computation, what's the difference between the below codes and log2_up/down codes?

    @@ -607,6 +678,26 @@ template<>
    Returns
    The CUDA error.
    +
    + + +

    ◆ find_log2()

    + +
    +
    +
    +template<typename value_t >
    + + + + + + + + +
    CUTLASS_HOST_DEVICE value_t cutlass::find_log2 (value_t x)
    +
    +
    @@ -638,82 +729,10 @@ template<typename value_t >

    Greatest common divisor

    -
    - - -

    ◆ get_Coord_dhw()

    - -
    -
    - - - - - - - - -
    CUTLASS_HOST_DEVICE Coord<3> cutlass::get_Coord_dhw (Coord< 4 > const & coord)
    -
    - -
    -
    - -

    ◆ get_Coord_hw() [1/2]

    - -
    -
    - - - - - - - - -
    CUTLASS_HOST_DEVICE Coord<2> cutlass::get_Coord_hw (Coord< 3 > const & coord)
    -
    - -
    -
    - -

    ◆ get_Coord_hw() [2/2]

    - -
    -
    - - - - - - - - -
    CUTLASS_HOST_DEVICE Coord<2> cutlass::get_Coord_hw (Coord< 4 > const & coord)
    -
    - -
    -
    - -

    ◆ get_Coord_hwc()

    - -
    -
    - - - - - - - - -
    CUTLASS_HOST_DEVICE Coord<3> cutlass::get_Coord_hwc (Coord< 4 > const & coord)
    -
    -
    -

    ◆ iterator_load() [1/4]

    +

    ◆ iterator_load()

    @@ -740,238 +759,10 @@ template<typename InputIterator , typename Fragment >
    -
    - - -

    ◆ iterator_load() [2/4]

    - -
    -
    -
    -template<typename InputIterator , typename Fragment , typename ConstPredicateAdapter >
    - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - -
    CUTLASS_HOST_DEVICE void cutlass::iterator_load (InputIterator const & _iterator,
    Fragment & fragment,
    typename InputIterator::Index offset,
    ConstPredicateAdapter predicate_adapter 
    )
    -
    - -
    -
    - -

    ◆ iterator_load() [3/4]

    - -
    -
    -
    -template<typename InputIterator , typename Fragment >
    - - - - - - - - - - - - - - - - - - - - - - - - -
    CUTLASS_HOST_DEVICE void cutlass::iterator_load (InputIterator const & iterator,
    Fragment & fragment,
    typename InputIterator::Index offset = 0 
    )
    -
    - -
    -
    - -

    ◆ iterator_load() [4/4]

    - -
    -
    -
    -template<typename InputIterator , typename Fragment , typename ConstPredicateAdapter >
    - - - - - - - - - - - - - - - - - - - - - - - - -
    CUTLASS_HOST_DEVICE void cutlass::iterator_load (InputIterator const & iterator,
    Fragment & fragment,
    ConstPredicateAdapter pred_it 
    )
    -
    - -
    -
    - -

    ◆ iterator_load_post_increment() [1/3]

    - -
    -
    -
    -template<typename InputIterator , typename Fragment , typename ConstPredicateAdapter >
    - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - -
    CUTLASS_HOST_DEVICE void cutlass::iterator_load_post_increment (InputIterator & iterator,
    Fragment & fragment,
    typename InputIterator::Index offset,
    ConstPredicateAdapter predicate_adapter 
    )
    -
    - -
    -
    - -

    ◆ iterator_load_post_increment() [2/3]

    - -
    -
    -
    -template<typename InputIterator , typename Fragment >
    - - - - - - - - - - - - - - - - - - - - - - - - -
    CUTLASS_HOST_DEVICE void cutlass::iterator_load_post_increment (InputIterator & iterator,
    Fragment & fragment,
    typename InputIterator::Index offset = 0 
    )
    -
    - -
    -
    - -

    ◆ iterator_load_post_increment() [3/3]

    - -
    -
    -
    -template<typename InputIterator , typename Fragment , typename ConstPredicateAdapter >
    - - - - - - - - - - - - - - - - - - - - - - - - -
    CUTLASS_HOST_DEVICE void cutlass::iterator_load_post_increment (InputIterator & iterator,
    Fragment & fragment,
    ConstPredicateAdapter pred_it 
    )
    -
    -
    -

    ◆ iterator_store() [1/4]

    +

    ◆ iterator_store()

    @@ -998,234 +789,6 @@ template<typename OutputIterator , typename Fragment >
    -
    - - -

    ◆ iterator_store() [2/4]

    - -
    -
    -
    -template<typename OutputIterator , typename Fragment , typename ConstPredicateAdapter >
    - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - -
    CUTLASS_HOST_DEVICE void cutlass::iterator_store (OutputIterator const & _iterator,
    Fragment const & fragment,
    typename OutputIterator::Index offset,
    ConstPredicateAdapter predicate_adapter 
    )
    -
    - -
    -
    - -

    ◆ iterator_store() [3/4]

    - -
    -
    -
    -template<typename OutputIterator , typename Fragment >
    - - - - - - - - - - - - - - - - - - - - - - - - -
    CUTLASS_HOST_DEVICE void cutlass::iterator_store (OutputIterator const & iterator,
    Fragment const & fragment,
    typename OutputIterator::Index offset = 0 
    )
    -
    - -
    -
    - -

    ◆ iterator_store() [4/4]

    - -
    -
    -
    -template<typename OutputIterator , typename Fragment , typename ConstPredicateAdapter >
    - - - - - - - - - - - - - - - - - - - - - - - - -
    CUTLASS_HOST_DEVICE void cutlass::iterator_store (OutputIterator const & iterator,
    Fragment const & fragment,
    ConstPredicateAdapter pred_it 
    )
    -
    - -
    -
    - -

    ◆ iterator_store_post_increment() [1/3]

    - -
    -
    -
    -template<typename OutputIterator , typename Fragment , typename ConstPredicateAdapter >
    - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - -
    CUTLASS_HOST_DEVICE void cutlass::iterator_store_post_increment (OutputIterator & iterator,
    Fragment const & fragment,
    typename OutputIterator::Index offset,
    ConstPredicateAdapter predicate_adapter 
    )
    -
    - -
    -
    - -

    ◆ iterator_store_post_increment() [2/3]

    - -
    -
    -
    -template<typename OutputIterator , typename Fragment >
    - - - - - - - - - - - - - - - - - - - - - - - - -
    CUTLASS_HOST_DEVICE void cutlass::iterator_store_post_increment (OutputIterator & iterator,
    Fragment const & fragment,
    typename OutputIterator::Index offset = 0 
    )
    -
    - -
    -
    - -

    ◆ iterator_store_post_increment() [3/3]

    - -
    -
    -
    -template<typename OutputIterator , typename Fragment , typename ConstPredicateAdapter >
    - - - - - - - - - - - - - - - - - - - - - - - - -
    CUTLASS_HOST_DEVICE void cutlass::iterator_store_post_increment (OutputIterator & iterator,
    Fragment const & fragment,
    ConstPredicateAdapter pred_it 
    )
    -
    -
    @@ -1379,8 +942,27 @@ template<typename value_t > - -

    ◆ make_zero() [1/2]

    + +

    ◆ make_Coord_from_shape()

    + +
    +
    +
    +template<typename Shape_ >
    + + + + + + + +
    CUTLASS_HOST_DEVICE Coord<3> cutlass::make_Coord_from_shape ()
    +
    + +
    +
    + +

    ◆ make_zero() [1/2]

    @@ -1388,7 +970,7 @@ template<typename value_t >
    template<typename Scalar_ >
    - + @@ -1399,8 +981,8 @@ template<typename Scalar_ > - -

    ◆ make_zero() [2/2]

    + +

    ◆ make_zero() [2/2]

    @@ -1408,7 +990,7 @@ template<typename Scalar_ >
    template<typename Scalar_ , int kLanes_>
    CUTLASS_DEVICE void cutlass::make_zero CUTLASS_HOST_DEVICE void cutlass::make_zero ( Scalar_ &  x)
    - + @@ -1417,6 +999,354 @@ template<typename Scalar_ , int kLanes_>
    CUTLASS_DEVICE void cutlass::make_zero CUTLASS_HOST_DEVICE void cutlass::make_zero ( Vector< Scalar_, kLanes_ > &  vec)
    +
    + + +

    ◆ make_ZipConvert()

    + +
    +
    +
    +template<typename First , typename Second >
    + + + + + + + + + + + + + + + + + + +
    CUTLASS_HOST_DEVICE ZipConvert<First, Second> cutlass::make_ZipConvert (First const & first,
    Second const & second 
    )
    +
    + +
    +
    + +

    ◆ make_ZipFragment()

    + +
    +
    +
    +template<typename First , typename Second >
    + + + + + + + + + + + + + + + + + + +
    CUTLASS_HOST_DEVICE ZipFragment<First, Second> cutlass::make_ZipFragment (First const & first,
    Second const & second 
    )
    +
    + +
    +
    + +

    ◆ make_ZipTensorRef()

    + +
    +
    +
    +template<typename First , typename Second >
    + + + + + + + + + + + + + + + + + + +
    CUTLASS_HOST_DEVICE ZipTensorRef<First, Second> cutlass::make_ZipTensorRef (First const & first,
    Second const & second 
    )
    +
    + +
    +
    + +

    ◆ operator<<() [1/7]

    + +
    +
    +
    +template<int Rank>
    + + + + + + + + + + + + + + + + + + +
    std::ostream& cutlass::operator<< (std::ostream & out,
    Coord< Rank > const & coord 
    )
    +
    + +
    +
    + +

    ◆ operator<<() [2/7]

    + +
    +
    +
    +template<typename T >
    + + + + + +
    + + + + + + + + + + + + + + + + + + +
    std::ostream& cutlass::operator<< (std::ostream & out,
    ScalarIO< T > const & scalar 
    )
    +     		+inline
    +
    + +
    +
    + +

    ◆ operator<<() [3/7]

    + +
    +
    +
    +template<>
    + + + + + +
    + + + + + + + + + + + + + + + + + + +
    std::ostream& cutlass::operator<< (std::ostream & out,
    ScalarIO< int8_t > const & scalar 
    )
    +     		+inline
    +
    + +
    +
    + +

    ◆ operator<<() [4/7]

    + +
    +
    +
    +template<>
    + + + + + +
    + + + + + + + + + + + + + + + + + + +
    std::ostream& cutlass::operator<< (std::ostream & out,
    ScalarIO< uint8_t > const & scalar 
    )
    +     		+inline
    +
    + +
    +
    + +

    ◆ operator<<() [5/7]

    + +
    +
    +
    +template<>
    + + + + + +
    + + + + + + + + + + + + + + + + + + +
    std::ostream& cutlass::operator<< (std::ostream & out,
    ScalarIO< cutlass::Vector< cutlass::bin1_t, 32 > > const & scalar 
    )
    +     		+inline
    +
    + +
    +
    + +

    ◆ operator<<() [6/7]

    + +
    +
    +
    +template<>
    + + + + + +
    + + + + + + + + + + + + + + + + + + +
    std::ostream& cutlass::operator<< (std::ostream & out,
    ScalarIO< cutlass::Vector< cutlass::int4_t, 8 > > const & scalar 
    )
    +     		+inline
    +
    + +
    +
    + +

    ◆ operator<<() [7/7]

    + +
    +
    +
    +template<>
    + + + + + +
    + + + + + + + + + + + + + + + + + + +
    std::ostream& cutlass::operator<< (std::ostream & out,
    ScalarIO< cutlass::Vector< cutlass::uint4_t, 8 > > const & scalar 
    )
    +     		+inline
    +
    +
    @@ -1448,108 +1378,12 @@ template<typename dividend_t , typename divisor_t >

    Round dividend up to the nearest multiple of divisor

    -
    - - -

    ◆ shared_iterator_load() [1/2]

    - -
    -
    -
    -template<typename InputIterator , typename Fragment >
    - - - - - - - - - - - - - - - - - - -
    CUTLASS_DEVICE void cutlass::shared_iterator_load (InputIterator & iterator,
    Fragment & fragment 
    )
    -
    - -
    -
    - -

    ◆ shared_iterator_load() [2/2]

    - -
    -
    -
    -template<typename InputIterator , typename Fragment >
    - - - - - - - - - - - - - - - - - - - - - - - - -
    CUTLASS_DEVICE void cutlass::shared_iterator_load (InputIterator & iterator,
    Fragment & fragment,
    int d 
    )
    -
    - -
    -
    - -

    ◆ shared_iterator_store()

    - -
    -
    -
    -template<typename OutputIterator , typename Fragment >
    - - - - - - - - - - - - - - - - - - -
    CUTLASS_DEVICE void cutlass::shared_iterator_store (OutputIterator & iterator,
    Fragment const & fragment 
    )
    -
    -
    -Generated on Thu May 3 2018 16:36:12 for Cutlass by   +Generated on Wed Sep 19 2018 13:58:50 for Cutlass by   doxygen 1.8.14
    diff --git a/docs/namespacecutlass_1_1detail.html b/docs/namespacecutlass_1_1detail.html new file mode 100644 index 00000000..154ce5c4 --- /dev/null +++ b/docs/namespacecutlass_1_1detail.html @@ -0,0 +1,95 @@ + + + + + + + +Cutlass: cutlass::detail Namespace Reference + + + + + + + + + + +
    +
    + + + + + + +
    +
    Cutlass +
    +
    CUDA Templates for Linear Algebra Subroutines and Solvers
    +
    +
    + + + + + + + + +
    +
    + + +
    + +
    + + +
    +
    +
    +Classes
    +
    +
    cutlass::detail Namespace Reference
    +
    +
    + + + + +

    +Classes

    class  ScalarOrPointer
     
    +
    + +
    +Generated on Wed Sep 19 2018 13:58:51 for Cutlass by   +doxygen + 1.8.14 +
    + + diff --git a/docs/namespacecutlass_1_1gemm.html b/docs/namespacecutlass_1_1gemm.html index 1c84e448..1545f43f 100644 --- a/docs/namespacecutlass_1_1gemm.html +++ b/docs/namespacecutlass_1_1gemm.html @@ -84,19 +84,28 @@ $(function() { Classes struct  ClearAccumulators   +struct  ColumnMajorBlockSwizzle +  struct  DgemmConfig   struct  DgemmTraits   +struct  Fp16SgemmConfig +  +struct  Fp16SgemmSgemmTraits +  struct  FragmentMultiplyAdd   -struct  FragmentMultiplyAdd< half > +struct  FragmentMultiplyAdd< half, half, true >   struct  Gemm   struct  GemmConfig   +struct  GemmCoord +  struct  GemmDesc + GEMM problem description. More...
      struct  GemmEpilogue   @@ -151,7 +160,8 @@ Classes   struct  GlobalLoadStream   -struct  GlobalLoadStreamBase +struct  GlobalLoadStreamPair + Collect the global load streams for multiplicands. More...
      struct  HgemmConfig   @@ -187,9 +197,7 @@ Classes   struct  IgemmConfig   -struct  IgemmConfig< OutputTile_, int8_t, AccumulatorsPerThread_ > -  -struct  IgemmContiguousGlobalTileTraits +struct  IgemmConfig< OutputTile_, int8_t, ThreadGemmShape_ >   struct  IgemmEpilogue   @@ -205,6 +213,8 @@ Classes   struct  IgemmFloatToInt8Converter   +struct  IgemmGlobalIteratorAb +  struct  IgemmGlobalLoadTransformer   struct  IgemmGlobalLoadTransformer< Fragment< int8_t, kElements_ >, float > @@ -213,6 +223,8 @@ Classes   struct  IgemmGlobalStoreTransformer< float, Fragment< int8_t, kElements_ > >   +struct  IgemmGlobalTileTraits +  struct  IgemmInt8ToFloatConverter   struct  IgemmSharedStoreTransformer @@ -221,11 +233,15 @@ Classes   struct  IgemmTileTraitsHelperA   -struct  IgemmTileTraitsHelperA< MatrixLayout::kColumnMajor, GemmConfig_ > +struct  IgemmTileTraitsHelperA< MatrixLayout::kColumnMajor, GemmConfig_, Index_ > +  +struct  IgemmTileTraitsHelperA< MatrixLayout::kRowMajor, GemmConfig_, Index_ >   struct  IgemmTileTraitsHelperB   -struct  IgemmTileTraitsHelperB< MatrixLayout::kRowMajor, GemmConfig_ > +struct  IgemmTileTraitsHelperB< MatrixLayout::kColumnMajor, GemmConfig_, Index_ > +  +struct  IgemmTileTraitsHelperB< MatrixLayout::kRowMajor, GemmConfig_, Index_ >   struct  IgemmTraits   @@ -243,9 +259,17 @@ Classes   struct  IgemmTransformerB< MatrixLayout::kRowMajor, Iterator_ >   +struct  Launch + Partial specialization for launching the GEMM kernel with or without launch bounds. More...
    +  +struct  Launch< Gemm, false > + Partial specialization for launching the GEMM kernel with or without launch bounds. More...
    +  struct  LinearScaling  Functor to compute linear combination of fragments. More...
      +struct  LinearScalingDevicePtr +  struct  ProjectOperand   struct  ProjectOperand< GemmOperand::kA, Kstrided > @@ -264,26 +288,39 @@ Classes   struct  ReshapeThreads< Tile_, Threads_, true >   +struct  RowMajorBlockSwizzle +  struct  SgemmConfig   +struct  SgemmLBTraits + Helper to define SGEMM traits using Launch Bounds. More...
    +  struct  SgemmTraits   struct  SharedLoadStream   +struct  SharedStreamPair + Collect the global load streams for multiplicands. More...
    +  struct  SimplifiedGemmEpilogueTraits   struct  SimplifiedGemmTraits   struct  SimplifiedGemmTraitsHelper   +struct  swizzleDirection +  struct  ThreadMultiplyAdd  Template performing matrix multiply-add operation within a thread. More...
      -struct  ThreadMultiplyAdd< AccumulatorsPerThread_, ThreadsPerWarp_, half, half, half > - Template performing matrix multiply-add operation within a thread. More...
    +struct  ThreadMultiplyAdd< ThreadGemmShape_, ThreadsPerWarp_, half, half, float > + Template performing matrix multiply-add operation within a thread. More...
      -struct  ThreadMultiplyAdd< AccumulatorsPerThread_, ThreadsPerWarp_, int8_t, int8_t, int > - Template performing matrix multiply-add operation within a thread. More...
    +struct  ThreadMultiplyAdd< ThreadGemmShape_, ThreadsPerWarp_, half, half, half > + Template performing matrix multiply-add operation within a thread. More...
    +  +struct  ThreadMultiplyAdd< ThreadGemmShape_, ThreadsPerWarp_, int8_t, int8_t, int > + Template performing matrix multiply-add operation within a thread. More...
      struct  WmmaGemmGlobalIteratorCd   @@ -292,18 +329,29 @@ Classes - - - + + + + + + + + + + + + + +

    Functions

    template<typename Gemm_ >
    __global__ void gemm_kernel (typename Gemm_::Params params)
     
    template<typename Gemm_ >
    __global__ __launch_bounds__ (Gemm_::kThreads) void gemm_kernel(typename Gemm_
     GEMM kernel with launch bounds specified. More...
     
    template<typename Gemm_ >
    __global__ void gemm_kernel_nolb (typename Gemm_::Params params)
     GEMM kernel without launch bounds specified. More...
     
    template<typename T >
    CUTLASS_DEVICE bool is_zero (T x)
     
    CUTLASS_DEVICE bool is_zero (half x)
     
    template<enum swizzleDirection::Kind >
    CUTLASS_DEVICE int getLinearIdx (int groups)
     
    template<>
    CUTLASS_DEVICE int getLinearIdx< swizzleDirection::Boustrophedon > (int groups)
     

    Function Documentation

    - -

    ◆ gemm_kernel()

    + +

    ◆ __launch_bounds__()

    @@ -311,7 +359,27 @@ Functions template<typename Gemm_ >
    - + + + + + + +
    __global__ void cutlass::gemm::gemm_kernel __global__ cutlass::gemm::__launch_bounds__ (Gemm_::kThreads )
    +
    + +
    + + +

    ◆ gemm_kernel_nolb()

    + +
    +
    +
    +template<typename Gemm_ >
    + + + @@ -320,6 +388,46 @@ template<typename Gemm_ >
    __global__ void cutlass::gemm::gemm_kernel_nolb ( typename Gemm_::Params  params)
    +
    +
    + +

    ◆ getLinearIdx()

    + +
    +
    +
    +template<enum swizzleDirection::Kind >
    + + + + + + + + +
    CUTLASS_DEVICE int cutlass::gemm::getLinearIdx (int groups)
    +
    + +
    +
    + +

    ◆ getLinearIdx< swizzleDirection::Boustrophedon >()

    + +
    +
    +
    +template<>
    + + + + + + + + +
    CUTLASS_DEVICE int cutlass::gemm::getLinearIdx< swizzleDirection::Boustrophedon > (int groups)
    +
    +
    @@ -363,7 +471,7 @@ template<typename T >
    -Generated on Thu May 3 2018 16:36:13 for Cutlass by   +Generated on Wed Sep 19 2018 13:58:51 for Cutlass by   doxygen 1.8.14
    diff --git a/docs/namespacecutlass_1_1platform.html b/docs/namespacecutlass_1_1platform.html index 2bf30c0d..b62a896a 100644 --- a/docs/namespacecutlass_1_1platform.html +++ b/docs/namespacecutlass_1_1platform.html @@ -122,6 +122,8 @@ Classes struct  bool_constant  std::bool_constant More...
      +class  complex +  struct  conditional  std::conditional (true specialization) More...
      @@ -256,6 +258,157 @@ Typedefs + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + @@ -399,8 +552,8 @@ template<> - -

    ◆ __align__() [3/13]

    + +

    ◆ __align__() [3/13]

    @@ -410,7 +563,7 @@ template<>
    - + @@ -419,8 +572,8 @@ template<> - -

    ◆ __align__() [4/13]

    + +

    ◆ __align__() [4/13]

    @@ -430,7 +583,7 @@ template<>
    - + @@ -439,8 +592,8 @@ template<> - -

    ◆ __align__() [5/13]

    + +

    ◆ __align__() [5/13]

    @@ -450,7 +603,7 @@ template<>
    - + @@ -477,10 +630,90 @@ template<>

    Functions

    CUTLASS_HOST_DEVICE float const & real (cuFloatComplex const &z)
     Returns the real part of the complex number. More...
     
    CUTLASS_HOST_DEVICE float & real (cuFloatComplex &z)
     Returns the real part of the complex number. More...
     
    CUTLASS_HOST_DEVICE double const & real (cuDoubleComplex const &z)
     Returns the real part of the complex number. More...
     
    CUTLASS_HOST_DEVICE double & real (cuDoubleComplex &z)
     Returns the real part of the complex number. More...
     
    CUTLASS_HOST_DEVICE float const & imag (cuFloatComplex const &z)
     Returns the imaginary part of the complex number. More...
     
    CUTLASS_HOST_DEVICE float & imag (cuFloatComplex &z)
     Returns the imaginary part of the complex number. More...
     
    CUTLASS_HOST_DEVICE double const & imag (cuDoubleComplex const &z)
     Returns the imaginary part of the complex number. More...
     
    CUTLASS_HOST_DEVICE double & imag (cuDoubleComplex &z)
     Returns the imaginary part of the complex number. More...
     
    template<typename T >
    CUTLASS_HOST_DEVICE T const & real (complex< T > const &z)
     Returns the real part of the complex number. More...
     
    template<typename T >
    CUTLASS_HOST_DEVICE T & real (complex< T > &z)
     Returns the real part of the complex number. More...
     
    template<typename T >
    CUTLASS_HOST_DEVICE T const & imag (complex< T > const &z)
     Returns the imaginary part of the complex number. More...
     
    template<typename T >
    CUTLASS_HOST_DEVICE T & imag (complex< T > &z)
     Returns the imaginary part of the complex number. More...
     
    template<typename T >
    std::ostream & operator<< (std::ostream &out, complex< T > const &z)
     
    template<typename T >
    CUTLASS_HOST_DEVICE bool operator== (complex< T > const &lhs, complex< T > const &rhs)
     Equality operator. More...
     
    template<typename T >
    CUTLASS_HOST_DEVICE bool operator!= (complex< T > const &lhs, complex< T > const &rhs)
     Inequality operator. More...
     
    template<typename T >
    CUTLASS_HOST_DEVICE complex< T > operator+ (complex< T > const &lhs, complex< T > const &rhs)
     Addition. More...
     
    template<typename T >
    CUTLASS_HOST_DEVICE complex< T > operator- (complex< T > const &lhs, complex< T > const &rhs)
     Subtraction. More...
     
    template<typename T >
    CUTLASS_HOST_DEVICE complex< T > operator* (complex< T > const &lhs, complex< T > const &rhs)
     Multiplication. More...
     
    template<typename T >
    CUTLASS_HOST_DEVICE complex< T > operator* (complex< T > const &lhs, T const &s)
     Scalar Multiplication. More...
     
    template<typename T >
    CUTLASS_HOST_DEVICE complex< T > operator* (T const &s, complex< T > const &rhs)
     Scalar Multiplication. More...
     
    template<typename T >
    CUTLASS_HOST_DEVICE complex< T > operator/ (complex< T > const &lhs, complex< T > const &rhs)
     Division. More...
     
    template<typename T >
    CUTLASS_HOST_DEVICE complex< T > operator/ (complex< T > const &lhs, T const &s)
     Scalar Division. More...
     
    template<typename T >
    CUTLASS_HOST_DEVICE complex< T > operator/ (T const &s, complex< T > const &rhs)
     Scalar divided by complex. More...
     
    template<typename T >
    CUTLASS_HOST_DEVICE complex< T > & operator+= (complex< T > &lhs, complex< T > const &rhs)
     Addition. More...
     
    template<typename T >
    CUTLASS_HOST_DEVICE complex< T > & operator-= (complex< T > &lhs, complex< T > const &rhs)
     Subtraction. More...
     
    template<typename T >
    CUTLASS_HOST_DEVICE complex< T > & operator*= (complex< T > &lhs, complex< T > const &rhs)
     Multiplication. More...
     
    template<typename T >
    CUTLASS_HOST_DEVICE complex< T > & operator*= (complex< T > &lhs, T s)
     Scalar multiplication. More...
     
    template<typename T >
    CUTLASS_HOST_DEVICE complex< T > & operator/= (complex< T > &lhs, complex< T > const &rhs)
     Division. More...
     
    template<typename T >
    CUTLASS_HOST_DEVICE T abs (complex< T > const &z)
     Returns the magnitude of the complex number. More...
     
    template<typename T >
    CUTLASS_HOST_DEVICE T arg (complex< T > const &z)
     Returns the magnitude of the complex number. More...
     
    template<typename T >
    CUTLASS_HOST_DEVICE T norm (complex< T > const &z)
     Returns the squared magnitude. More...
     
    template<typename T >
    CUTLASS_HOST_DEVICE complex< T > conj (complex< T > const &z)
     Returns the complex conjugate. More...
     
    template<typename T >
    CUTLASS_HOST_DEVICE complex< T > proj (complex< T > const &z)
     Projects the complex number z onto the Riemann sphere. More...
     
    template<typename T >
    CUTLASS_HOST_DEVICE complex< T > polar (T const &r, T const &theta=T())
     Returns a complex number with magnitude r and phase theta. More...
     
    template<typename T >
    CUTLASS_HOST_DEVICE complex< T > exp (complex< T > const &z)
     Computes the complex exponential of z. More...
     
    template<typename T >
    CUTLASS_HOST_DEVICE complex< T > log (complex< T > const &z)
     Computes the complex exponential of z. More...
     
    template<typename T >
    CUTLASS_HOST_DEVICE complex< T > log10 (complex< T > const &z)
     Computes the complex exponential of z. More...
     
    template<typename T >
    CUTLASS_HOST_DEVICE complex< T > sqrt (complex< T > const &z)
     Computes the square root of complex number z. More...
     
    template<typename T >
    CUTLASS_HOST_DEVICE complex< T > cos (complex< T > const &z)
     Computes the cosine of complex z. More...
     
    template<typename T >
    CUTLASS_HOST_DEVICE complex< T > sin (complex< T > const &z)
     Computes the sin of complex z. More...
     
    template<typename T >
    CUTLASS_HOST_DEVICE constexpr const T & min (const T &a, const T &b)
     std::min More...
    struct cutlass::platform::__align__ (64 4  )
    struct cutlass::platform::__align__ (128 8  )
    struct cutlass::platform::__align__ (256 16  )
    +
    + + +

    ◆ __align__() [7/13]

    + +
    +
    +
    +template<>
    + + + + + + + + +
    struct cutlass::platform::__align__ (32 )
    +
    + +
    +
    + +

    ◆ __align__() [8/13]

    + +
    +
    +
    +template<>
    + + + + + + + + +
    struct cutlass::platform::__align__ (64 )
    +
    + +
    +
    + +

    ◆ __align__() [9/13]

    + +
    +
    +
    +template<>
    + + + + + + + + +
    struct cutlass::platform::__align__ (128 )
    +
    + +
    +
    + +

    ◆ __align__() [10/13]

    + +
    +
    +
    +template<>
    + + + + + + + + +
    struct cutlass::platform::__align__ (256 )
    +
    +
    -

    ◆ __align__() [7/13]

    +

    ◆ __align__() [11/13]

    @@ -500,7 +733,7 @@ template<>
    -

    ◆ __align__() [8/13]

    +

    ◆ __align__() [12/13]

    @@ -520,7 +753,7 @@ template<>
    -

    ◆ __align__() [9/13]

    +

    ◆ __align__() [13/13]

    @@ -539,19 +772,19 @@ template<>
    - -

    ◆ __align__() [10/13]

    + +

    ◆ abs()

    -template<>
    +template<typename T >
    - + - - + +
    struct cutlass::platform::__align__ CUTLASS_HOST_DEVICE T cutlass::platform::abs (32 )complex< T > const & z)
    @@ -559,19 +792,19 @@ template<>
    - -

    ◆ __align__() [11/13]

    + +

    ◆ arg()

    -template<>
    +template<typename T >
    - + - - + +
    struct cutlass::platform::__align__ CUTLASS_HOST_DEVICE T cutlass::platform::arg (4 )complex< T > const & z)
    @@ -579,19 +812,19 @@ template<>
    - -

    ◆ __align__() [12/13]

    + +

    ◆ conj()

    -template<>
    +template<typename T >
    - + - - + +
    struct cutlass::platform::__align__ CUTLASS_HOST_DEVICE complex<T> cutlass::platform::conj (8 )complex< T > const & z)
    @@ -599,19 +832,191 @@ template<>
    - -

    ◆ __align__() [13/13]

    + +

    ◆ cos()

    -template<>
    +template<typename T >
    - + - - + + + + +
    struct cutlass::platform::__align__ CUTLASS_HOST_DEVICE complex<T> cutlass::platform::cos (16 )complex< T > const & z)
    +
    + +
    + + +

    ◆ exp()

    + +
    +
    +
    +template<typename T >
    + + + + + + + + +
    CUTLASS_HOST_DEVICE complex<T> cutlass::platform::exp (complex< T > const & z)
    +
    + +
    +
    + +

    ◆ imag() [1/6]

    + +
    +
    + + + + + + + + +
    CUTLASS_HOST_DEVICE float const& cutlass::platform::imag (cuFloatComplex const & z)
    +
    + +
    +
    + +

    ◆ imag() [2/6]

    + +
    +
    + + + + + + + + +
    CUTLASS_HOST_DEVICE float& cutlass::platform::imag (cuFloatComplex & z)
    +
    + +
    +
    + +

    ◆ imag() [3/6]

    + +
    +
    + + + + + + + + +
    CUTLASS_HOST_DEVICE double const& cutlass::platform::imag (cuDoubleComplex const & z)
    +
    + +
    +
    + +

    ◆ imag() [4/6]

    + +
    +
    + + + + + + + + +
    CUTLASS_HOST_DEVICE double& cutlass::platform::imag (cuDoubleComplex & z)
    +
    + +
    +
    + +

    ◆ imag() [5/6]

    + +
    +
    +
    +template<typename T >
    + + + + + + + + +
    CUTLASS_HOST_DEVICE T const& cutlass::platform::imag (complex< T > const & z)
    +
    + +
    +
    + +

    ◆ imag() [6/6]

    + +
    +
    +
    +template<typename T >
    + + + + + + + + +
    CUTLASS_HOST_DEVICE T& cutlass::platform::imag (complex< T > & z)
    +
    + +
    +
    + +

    ◆ log()

    + +
    +
    +
    +template<typename T >
    + + + + + + + + +
    CUTLASS_HOST_DEVICE complex<T> cutlass::platform::log (complex< T > const & z)
    +
    + +
    +
    + +

    ◆ log10()

    + +
    +
    +
    +template<typename T >
    + + + + + +
    CUTLASS_HOST_DEVICE complex<T> cutlass::platform::log10 (complex< T > const & z)
    @@ -707,10 +1112,30 @@ template<typename T >
    +
    + + +

    ◆ norm()

    + +
    +
    +
    +template<typename T >
    + + + + + + + + +
    CUTLASS_HOST_DEVICE T cutlass::platform::norm (complex< T > const & z)
    +
    +
    -

    ◆ operator!=()

    +

    ◆ operator!=() [1/2]

    @@ -737,6 +1162,426 @@ template<class T1 , class T2 >
    +
    + + +

    ◆ operator!=() [2/2]

    + +
    +
    +
    +template<typename T >
    + + + + + + + + + + + + + + + + + + +
    CUTLASS_HOST_DEVICE bool cutlass::platform::operator!= (complex< T > const & lhs,
    complex< T > const & rhs 
    )
    +
    + +
    +
    + +

    ◆ operator*() [1/3]

    + +
    +
    +
    +template<typename T >
    + + + + + + + + + + + + + + + + + + +
    CUTLASS_HOST_DEVICE complex<T> cutlass::platform::operator* (complex< T > const & lhs,
    complex< T > const & rhs 
    )
    +
    + +
    +
    + +

    ◆ operator*() [2/3]

    + +
    +
    +
    +template<typename T >
    + + + + + + + + + + + + + + + + + + +
    CUTLASS_HOST_DEVICE complex<T> cutlass::platform::operator* (complex< T > const & lhs,
    T const & s 
    )
    +
    + +
    +
    + +

    ◆ operator*() [3/3]

    + +
    +
    +
    +template<typename T >
    + + + + + + + + + + + + + + + + + + +
    CUTLASS_HOST_DEVICE complex<T> cutlass::platform::operator* (T const & s,
    complex< T > const & rhs 
    )
    +
    + +
    +
    + +

    ◆ operator*=() [1/2]

    + +
    +
    +
    +template<typename T >
    + + + + + + + + + + + + + + + + + + +
    CUTLASS_HOST_DEVICE complex<T>& cutlass::platform::operator*= (complex< T > & lhs,
    complex< T > const & rhs 
    )
    +
    + +
    +
    + +

    ◆ operator*=() [2/2]

    + +
    +
    +
    +template<typename T >
    + + + + + + + + + + + + + + + + + + +
    CUTLASS_HOST_DEVICE complex<T>& cutlass::platform::operator*= (complex< T > & lhs,
    T s 
    )
    +
    + +
    +
    + +

    ◆ operator+()

    + +
    +
    +
    +template<typename T >
    + + + + + + + + + + + + + + + + + + +
    CUTLASS_HOST_DEVICE complex<T> cutlass::platform::operator+ (complex< T > const & lhs,
    complex< T > const & rhs 
    )
    +
    + +
    +
    + +

    ◆ operator+=()

    + +
    +
    +
    +template<typename T >
    + + + + + + + + + + + + + + + + + + +
    CUTLASS_HOST_DEVICE complex<T>& cutlass::platform::operator+= (complex< T > & lhs,
    complex< T > const & rhs 
    )
    +
    + +
    +
    + +

    ◆ operator-()

    + +
    +
    +
    +template<typename T >
    + + + + + + + + + + + + + + + + + + +
    CUTLASS_HOST_DEVICE complex<T> cutlass::platform::operator- (complex< T > const & lhs,
    complex< T > const & rhs 
    )
    +
    + +
    +
    + +

    ◆ operator-=()

    + +
    +
    +
    +template<typename T >
    + + + + + + + + + + + + + + + + + + +
    CUTLASS_HOST_DEVICE complex<T>& cutlass::platform::operator-= (complex< T > & lhs,
    complex< T > const & rhs 
    )
    +
    + +
    +
    + +

    ◆ operator/() [1/3]

    + +
    +
    +
    +template<typename T >
    + + + + + + + + + + + + + + + + + + +
    CUTLASS_HOST_DEVICE complex<T> cutlass::platform::operator/ (complex< T > const & lhs,
    complex< T > const & rhs 
    )
    +
    + +
    +
    + +

    ◆ operator/() [2/3]

    + +
    +
    +
    +template<typename T >
    + + + + + + + + + + + + + + + + + + +
    CUTLASS_HOST_DEVICE complex<T> cutlass::platform::operator/ (complex< T > const & lhs,
    T const & s 
    )
    +
    + +
    +
    + +

    ◆ operator/() [3/3]

    + +
    +
    +
    +template<typename T >
    + + + + + + + + + + + + + + + + + + +
    CUTLASS_HOST_DEVICE complex<T> cutlass::platform::operator/ (T const & s,
    complex< T > const & rhs 
    )
    +
    + +
    +
    + +

    ◆ operator/=()

    + +
    +
    +
    +template<typename T >
    + + + + + + + + + + + + + + + + + + +
    CUTLASS_HOST_DEVICE complex<T>& cutlass::platform::operator/= (complex< T > & lhs,
    complex< T > const & rhs 
    )
    +
    +
    @@ -767,6 +1612,36 @@ template<class T1 , class T2 >
    +
    + + +

    ◆ operator<<()

    + +
    +
    +
    +template<typename T >
    + + + + + + + + + + + + + + + + + + +
    std::ostream& cutlass::platform::operator<< (std::ostream & out,
    complex< T > const & z 
    )
    +
    +
    @@ -797,10 +1672,40 @@ template<class T1 , class T2 >
    +
    + + +

    ◆ operator==() [1/2]

    + +
    +
    +
    +template<typename T >
    + + + + + + + + + + + + + + + + + + +
    CUTLASS_HOST_DEVICE bool cutlass::platform::operator== (complex< T > const & lhs,
    complex< T > const & rhs 
    )
    +
    +
    -

    ◆ operator==()

    +

    ◆ operator==() [2/2]

    @@ -887,6 +1792,208 @@ template<class T1 , class T2 >
    +
    + + +

    ◆ polar()

    + +
    +
    +
    +template<typename T >
    + + + + + + + + + + + + + + + + + + +
    CUTLASS_HOST_DEVICE complex<T> cutlass::platform::polar (T const & r,
    T const & theta = T() 
    )
    +
    + +
    +
    + +

    ◆ proj()

    + +
    +
    +
    +template<typename T >
    + + + + + + + + +
    CUTLASS_HOST_DEVICE complex<T> cutlass::platform::proj (complex< T > const & z)
    +
    + +
    +
    + +

    ◆ real() [1/6]

    + +
    +
    + + + + + + + + +
    CUTLASS_HOST_DEVICE float const& cutlass::platform::real (cuFloatComplex const & z)
    +
    + +
    +
    + +

    ◆ real() [2/6]

    + +
    +
    + + + + + + + + +
    CUTLASS_HOST_DEVICE float& cutlass::platform::real (cuFloatComplex & z)
    +
    + +
    +
    + +

    ◆ real() [3/6]

    + +
    +
    + + + + + + + + +
    CUTLASS_HOST_DEVICE double const& cutlass::platform::real (cuDoubleComplex const & z)
    +
    + +
    +
    + +

    ◆ real() [4/6]

    + +
    +
    + + + + + + + + +
    CUTLASS_HOST_DEVICE double& cutlass::platform::real (cuDoubleComplex & z)
    +
    + +
    +
    + +

    ◆ real() [5/6]

    + +
    +
    +
    +template<typename T >
    + + + + + + + + +
    CUTLASS_HOST_DEVICE T const& cutlass::platform::real (complex< T > const & z)
    +
    + +
    +
    + +

    ◆ real() [6/6]

    + +
    +
    +
    +template<typename T >
    + + + + + + + + +
    CUTLASS_HOST_DEVICE T& cutlass::platform::real (complex< T > & z)
    +
    + +
    +
    + +

    ◆ sin()

    + +
    +
    +
    +template<typename T >
    + + + + + + + + +
    CUTLASS_HOST_DEVICE complex<T> cutlass::platform::sin (complex< T > const & z)
    +
    + +
    +
    + +

    ◆ sqrt()

    + +
    +
    +
    +template<typename T >
    + + + + + + + + +
    CUTLASS_HOST_DEVICE complex<T> cutlass::platform::sqrt (complex< T > const & z)
    +
    +
    @@ -930,7 +2037,7 @@ template<typename T , typename Deleter >
    -Generated on Thu May 3 2018 16:36:14 for Cutlass by   +Generated on Wed Sep 19 2018 13:58:53 for Cutlass by   doxygen 1.8.14
    diff --git a/docs/namespacemembers.html b/docs/namespacemembers.html index 9566721d..a522eab7 100644 --- a/docs/namespacemembers.html +++ b/docs/namespacemembers.html @@ -73,22 +73,54 @@ $(function() {

    - _ -

    + + +

    - a -

    - c -

    +

    - e -

    + +

    - f -

    @@ -96,36 +128,30 @@ $(function() {
  • gcd() : cutlass
    • -
  • gemm_kernel() -: cutlass::gemm +
  • gemm_kernel_nolb() +: cutlass::gemm
    • -
  • get_Coord_dhw() -: cutlass +
  • getLinearIdx() +: cutlass::gemm
    • -
  • get_Coord_hw() -: cutlass -
    • -
  • get_Coord_hwc() -: cutlass +
  • getLinearIdx< swizzleDirection::Boustrophedon >() +: cutlass::gemm

    • - i -

    @@ -134,6 +160,12 @@ $(function() {
  • lcm() : cutlass
    • +
  • log() +: cutlass::platform +
    • +
  • log10() +: cutlass::platform +
    • @@ -141,11 +173,23 @@ $(function() {
  • make_Coord() : cutlass
    • +
  • make_Coord_from_shape() +: cutlass +
  • make_pair() : cutlass::platform
  • make_zero() -: cutlass +: cutlass +
    • +
  • make_ZipConvert() +: cutlass +
    • +
  • make_ZipFragment() +: cutlass +
    • +
  • make_ZipTensorRef() +: cutlass
  • max() : cutlass::platform @@ -156,18 +200,53 @@ $(function() { +

    - n -

    + +

    - o -

    +

    - p -

    + +

    - r -

    +

    - n -

    + +

    - o -

    +

    - p -

    + +

    - r -