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

/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/develop/include/ck_tile/core/arch/mma/amdgcn_mma.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/amdgcn_mma.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/numeric/vector_type.hpp"
 #include "ck_tile/core/utility/ignore.hpp"
  
 namespace ck_tile::core::arch::mma {
  
 struct Unsupported;
  
 #if defined(__cpp_concepts) && __cpp_concepts >= 201907L
 template <typename MmaOp>
 concept MmaOpI = requires(MmaOp op) {
     // Requires an op context
     typename MmaOp::OpType;
  
     // Captures types for inputs / outputs to mma function
     typename MmaOp::AVecType;
     typename MmaOp::BVecType;
     typename MmaOp::CVecType;
  
     // Captures CK-specific layout properties
     { MmaOp::kAMBlock } -> std::convertible_to<unsigned int>;
     { MmaOp::kBNBlock } -> std::convertible_to<unsigned int>;
     { MmaOp::kAMLane } -> std::convertible_to<unsigned int>;
     { MmaOp::kBNLane } -> std::convertible_to<unsigned int>;
     { MmaOp::kABKLane } -> std::convertible_to<unsigned int>;
     { MmaOp::kABKPerLane } -> std::convertible_to<unsigned int>;
     { MmaOp::kCMLane } -> std::convertible_to<unsigned int>;
     { MmaOp::kCNLane } -> std::convertible_to<unsigned int>;
     { MmaOp::kCM0PerLane } -> std::convertible_to<unsigned int>;
     { MmaOp::kCM1PerLane } -> std::convertible_to<unsigned int>;
  
     // Static exec function
     {
         MmaOp::exec(
             typename MmaOp::AVecType{}, typename MmaOp::BVecType{}, typename MmaOp::CVecType{})
     } -> std::convertible_to<typename MmaOp::CVecType>;
 };
  
 #endif // defined(__cpp_concepts) && __cpp_concepts >= 201907L
  
 template <typename ADataType,
           typename BDataType,
           typename CDataType,
           uint32_t BlockM,
           uint32_t BlockN,
           uint32_t BlockK,
           typename CtrlFlags,
           typename CompilerTarget,
           typename Enabler = void>
 struct amdgcn_mma
 {
     // The base instance is unsupported because there is no __builtin to wrap.
     using OpType = Unsupported;
  
     // Interface types for A, B, C vectors types
     using AVecType = ext_vector_t<ADataType, 1>;
     using BVecType = ext_vector_t<BDataType, 1>;
     using CVecType = ext_vector_t<CDataType, 1>;
  
     // Layout constants - default to 0
     static constexpr index_t kAMBlock = 0;
     static constexpr index_t kBNBlock = 0;
  
     static constexpr index_t kAMLane     = 0;
     static constexpr index_t kBNLane     = 0;
     static constexpr index_t kABKLane    = 0;
     static constexpr index_t kABKPerLane = 0;
  
     static constexpr index_t kCMLane     = 0;
     static constexpr index_t kCNLane     = 0;
     static constexpr index_t kCM0PerLane = 0;
     static constexpr index_t kCM1PerLane = 0;
  
     // This is a default pass-through implementation that doesn't do anything practical.
     CK_TILE_DEVICE static CVecType const&
     exec(AVecType const& regsA, BVecType const& regsB, CVecType const& regsC)
     {
         ignore(regsA, regsB);
         return regsC; // No-op, just return C
     }
 };
  
 } // namespace ck_tile::core::arch::mma
  
 // Include the implementations
 #include "wmma/wmma.hpp"
 #include "mfma/mfma.hpp"