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

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

mma_transforms.hpp

Go to the documentation of this file.

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

2 // SPDX-License-Identifier: MIT

3 #pragma once

4 namespace ck_tile::core::arch::mma {

5 

10 struct PassThroughTransform

11 {

12  template <typename VecType>

13  CK_TILE_DEVICE static decltype(auto) exec(VecType&& v)

14  {

15  return std::forward<VecType>(v);

16  }

17 };

18 

26 template <typename MmaOp, typename CompilerTarget, typename Enable = void>

27 // TODO: c++20 template <MmaOpI MmaOp, amdgcn_target_arch_id CompilerTarget, typename Enable = void>

28 struct MmaTransformsDefaultSelector;

29 

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

31 

36 template <typename MmaTransforms>

37 concept MmaTransformsI = requires(MmaTransforms transforms) {

38  // Transforms should define TransformA, TransformB, TransformC, and TransformD types

39  typename MmaTransforms::ATransform;

40  typename MmaTransforms::BTransform;

41  typename MmaTransforms::CTransform;

42  typename MmaTransforms::DTransform;

43 };

44 

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

46 

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

CK_TILE_DEVICE

#define CK_TILE_DEVICE

Definition: config.hpp:45

ck_tile::core::arch::mma

Definition: amdgcn_mma.hpp:10

ck_tile::core::arch::mma::MmaTransformsDefaultSelector

Implements the default MMA transforms selection for gfx9 targets.

Definition: mma_transforms.hpp:28

ck_tile::core::arch::mma::PassThroughTransform

A no-op transform that passes through the input as-is.

Definition: mma_transforms.hpp:11

ck_tile::core::arch::mma::PassThroughTransform::exec

static decltype(auto) CK_TILE_DEVICE exec(VecType &&v)

Definition: mma_transforms.hpp:13