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

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

Composable Kernel: /home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/develop/include/ck_tile/core/arch/mma/wmma/wmma_transforms.hpp Source File

Go to the documentation of this file.

 // Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
 // SPDX-License-Identifier: MIT
  
 #pragma once
  
 #include "ck_tile/core/config.hpp"
 #include "ck_tile/core/arch/arch.hpp"
 #include "ck_tile/core/arch/mma/mma_transforms.hpp"
  
 namespace ck_tile::core::arch::mma {
  
 struct DuplicateTransform
 {
     template <typename VecType>
     CK_TILE_DEVICE static decltype(auto) exec(VecType&& v)
     {
         // TODO: Implement duplication logic to broadcast low
         // register elements to high elements [0 - (N/2 -1)] -> [N/2 - (N-1)]
         return std::forward<VecType>(v);
     }
 };
  
 struct PadTransform
 {
     template <typename VecType>
     CK_TILE_DEVICE static decltype(auto) exec(VecType&& v)
     {
         // TODO: Implement b32 padding logic.
         // E.g., for fp16, pad each 16-bit element with 16 bits of 0 to make 32-bit elements
         return std::forward<VecType>(v);
     }
 };
  
 struct UnpadTransform
 {
     template <typename VecType>
     CK_TILE_DEVICE static decltype(auto) exec(VecType&& v)
     {
         // TODO: Implement b32 logic to unpad 32 to original data type.
         return std::forward<VecType>(v);
     }
 };
  
 struct MmaDefaultTransformsGfx11
 {
     using ATransform = DuplicateTransform;
     using BTransform = DuplicateTransform;
     using CTransform = PadTransform;
     using DTransform = UnpadTransform;
 };
  
 struct MmaDefaultTransformsGfx12
 {
     using ATransform = PassThroughTransform;
     using BTransform = PassThroughTransform;
     using CTransform = PassThroughTransform;
     using DTransform = PassThroughTransform;
 };
  
 template <typename MmaOp, typename CompilerTarget>
 // TODO: c++20 template <MmaOpI MmaOp, amdgcn_target_arch_id GfxTargetId>
 // TODO: c++20 requires
 struct MmaTransformsDefaultSelector<MmaOp,
                                     CompilerTarget,
                                     enable_if_target_family_gfx11_t<CompilerTarget>>
 {
     using SelectedTransforms = MmaDefaultTransformsGfx11;
 };
  
 template <typename MmaOp, typename CompilerTarget>
 // TODO: c++20 template <MmaOpI MmaOp, amdgcn_target_arch_id GfxTargetId>
 // TODO: c++20 requires
 struct MmaTransformsDefaultSelector<MmaOp,
                                     CompilerTarget,
                                     enable_if_target_family_gfx12_t<CompilerTarget>>
 {
     using SelectedTransforms = MmaDefaultTransformsGfx12;
 };
  
 } // namespace ck_tile::core::arch::mma