Back to Cutlass

CUTLASS: mma_simt_policy.h Source File

docs/mma__simt__policy_8h_source.html

4.4.26.1 KB
Original Source

| | CUTLASS

CUDA Templates for Linear Algebra Subroutines and Solvers |

mma_simt_policy.h

[Go to the documentation of this file.](mma simt policy_8h.html)

1 /***************************************************************************************************

2 * Copyright (c) 2017-2019, NVIDIA CORPORATION. All rights reserved.

3 *

4 * Redistribution and use in source and binary forms, with or without modification, are permitted

5 * provided that the following conditions are met:

6 * * Redistributions of source code must retain the above copyright notice, this list of

7 * conditions and the following disclaimer.

8 * * Redistributions in binary form must reproduce the above copyright notice, this list of

9 * conditions and the following disclaimer in the documentation and/or other materials

10 * provided with the distribution.

11 * * Neither the name of the NVIDIA CORPORATION nor the names of its contributors may be used

12 * to endorse or promote products derived from this software without specific prior written

13 * permission.

14 *

15 * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND ANY EXPRESS OR

16 * IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND

17 * FITNESS FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE

18 * FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING,

19 * BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS;

20 * OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT,

21 * STRICT LIABILITY, OR TOR (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE

22 * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.

23 *

24 **************************************************************************************************/

30 #pragma once

31

32 #include "cutlass/cutlass.h"

33

34 namespace cutlass {

35 namespace gemm {

36 namespace warp {

37

39

41 template <

42typename WarpShape_,

43typename LaneLayout_,

44typename LaneMmaShape_

45 >

46 struct MmaSimtPolicy {

47using WarpShape = WarpShape_;

48using LaneLayout = LaneLayout_;

49using LaneMmaShape = LaneMmaShape_;

50using MmaShape = LaneMmaShape;

51

53CUTLASS_HOST_DEVICE

54static LaneLayout get_lane_layout() {

55return LaneLayout::packed({WarpShape::kRow, WarpShape::kColumn});

56 }

57 };

58

60

61 } // namespace warp

62 } // namespace gemm

63 } // namespace cutlass

cutlass::gemm::warp::MmaSimtPolicy::MmaShape

LaneMmaShape MmaShape

Definition: mma_simt_policy.h:50

cutlass

Definition: aligned_buffer.h:35

cutlass::gemm::warp::MmaSimtPolicy

Describes the arrangement and configuration of per-lane operations in warp-level matrix multiply...

Definition: mma_simt_policy.h:46

CUTLASS_HOST_DEVICE

#define CUTLASS_HOST_DEVICE

Definition: cutlass.h:89

cutlass::gemm::warp::MmaSimtPolicy::get_lane_layout

static CUTLASS_HOST_DEVICE LaneLayout get_lane_layout()

Returns a layout functor mapping lane position in the warp to thread ID.

Definition: mma_simt_policy.h:54

cutlass::gemm::warp::MmaSimtPolicy::WarpShape

WarpShape_ WarpShape

Definition: mma_simt_policy.h:47

cutlass::gemm::warp::MmaSimtPolicy::LaneLayout

LaneLayout_ LaneLayout

Definition: mma_simt_policy.h:48

cutlass.h

Basic include for CUTLASS.

cutlass::gemm::warp::MmaSimtPolicy::LaneMmaShape

LaneMmaShape_ LaneMmaShape

Definition: mma_simt_policy.h:49


Generated by 1.8.11