/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/develop/include/ck_tile/core/arch/mma/mma_selector.hpp Source File

• develop
• include
• ck_tile
• core
• arch
• mma

mma_selector.hpp

Go to the documentation of this file.

1 // Copyright (c) Advanced Micro Devices, Inc., or its affiliates.

2 // SPDX-License-Identifier: MIT

3 

4 #pragma once

5 namespace ck_tile::core::arch::mma {

6 

28 template <typename ADataType,

29  typename BDataType,

30  typename CDataType,

31  uint32_t FragM,

32  uint32_t FragN,

33  uint32_t FragK,

34  typename CompilerTarget,

35  typename Enable = void>

36 // TODO c++20 requires

37 struct MmaDefaultSelector

38 {

39  // By default, no selection is made, and we fall back to a pass-through unsupported

40  // implementation. This is because we do not have any knowledge of the target architecture here.

41  using SelectedOp =

42  amdgcn_mma<ADataType, BDataType, CDataType, FragM, FragN, FragK, void, amdgcn_target<>>;

43 };

44 

45 #if defined(__cpp_concepts) && __cpp_concepts >= 201907L

46 

51 template <typename MmaSelector>

52 concept MmaSelectorI = requires(MmaSelector op) {

53  // Selectors should have a resulting SelectedOp type

54  typename MmaSelector::SelectedOp;

55 };

56 

57 #endif // defined(__cpp_concepts) && __cpp_concepts >= 201907L

58 

59 } // namespace ck_tile::core::arch::mma

60 

61 // Include the implementations

62 #include "wmma/wmma_selector.hpp"

63 #include "mfma/mfma_selector.hpp"

mfma_selector.hpp

ck_tile::core::arch::mma

Definition: amdgcn_mma.hpp:10

uint32_t

unsigned int uint32_t

Definition: stdint.h:126

ck_tile::core::arch::mma::MmaDefaultSelector

Implements the gfx9 default MMA selector strategy for wave-wise MMA decomposition....

Definition: mma_selector.hpp:38

ck_tile::core::arch::mma::amdgcn_mma

This is the default MmaOp policy. Instances of this class are to be used as MmaOp policies....

Definition: amdgcn_mma.hpp:82

wmma_selector.hpp