/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/develop/include/ck_tile/core/arch/mma/wmma/wmma_gfx12.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_gfx12.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_gfx12.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 "wmma_traits.hpp"
  
 #include "ck_tile/core/config.hpp"
 #include "ck_tile/core/arch/arch.hpp"
 #include "ck_tile/core/arch/mma/amdgcn_mma.hpp"
 #include "ck_tile/core/numeric/vector_type.hpp"
  
 namespace ck_tile::core::arch::mma {
  
 // NOTE: At this point forward, we are specializing amdgcn_mma for each target id as needed.
 // This is because some built-ins are only available on certain target ids.
 // We can also do things, such add some padding specializations for when we need to use
 // smaller values of K that aren't directly supported by the built-ins.
 // For flexibility, it is recommended that for each backend wrapper it supports at least
 // one packed register for each input to be able to process smaller K values by padding.
  
 // TODO: c++20 template <CtrlFlagsGfx12I CtrlFlags, amdgcn_target CompilerTarget>
 // TODO: c++20 requires
 template <typename CtrlFlags, typename CompilerTarget>
 struct amdgcn_mma<fp16_t,
                   fp16_t,
                   fp32_t,
                   16u,
                   16u,
                   16u,
                   CtrlFlags,
                   CompilerTarget,
                   enable_if_target_family_gfx12_t<CompilerTarget>>
 {
     // Wmma operation type
     using OpType = WmmaOp;
  
     // Register types
     using AVecType = ext_vector_t<fp16_t, 8>;
     using BVecType = ext_vector_t<fp16_t, 8>;
     using CVecType = ext_vector_t<fp32_t, 8>;
  
     // Layout constants
     static constexpr index_t kAMBlock    = 1;
     static constexpr index_t kBNBlock    = 1;
     static constexpr index_t kAMLane     = 16;
     static constexpr index_t kBNLane     = 16;
     static constexpr index_t kABKLane    = 8;
     static constexpr index_t kABKPerLane = 8;
     static constexpr index_t kCMLane     = 2;
     static constexpr index_t kCNLane     = 2;
     static constexpr index_t kCM0PerLane = 4;
     static constexpr index_t kCM1PerLane = 1;
  
     CK_TILE_DEVICE static auto
     exec(AVecType const& aVec, BVecType const& bVec, CVecType const& cVec) -> CVecType
     {
         return {__builtin_amdgcn_wmma_f32_16x16x16_f16_w32_gfx12(aVec, bVec, cVec)};
     }
 };
  
 } // namespace ck_tile::core::arch::mma