/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/develop/include/ck_tile/core/arch/mma/wmma/wmma.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.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.hpp Source File
wmma.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 
6 namespace ck_tile::core::arch::mma {
7 
12 enum struct WmmaCtrlFlags : bool
13 {
14  // Only has an effect on gfx11 when the accumulator is 16-bit
15  // Determines which half of the 32-bit accum register to use
16  // Low = bits [15:0]
17  // High = bits[31:16]
18  LOW = false,
19  HIGH = true,
20 
21  // Only has an effect on gfx11 / 12 when the input is 8-bit int
22  // Signage indicator of inputs / accum
23  UNSIGNED = false,
24  SIGNED = true
25 };
26 
27 } // namespace ck_tile::core::arch::mma
28 
29 // Include the architecture-specific WMMA implementations and traits
30 #include "wmma_gfx11.hpp"
31 #include "wmma_gfx12.hpp"
32 #include "wmma_selector.hpp"
33 #include "wmma_traits.hpp"
34 #include "wmma_transforms.hpp"
Definition: amdgcn_mma.hpp:10
WmmaCtrlFlags
Common wmma control flags for gfx11 and gfx12.
Definition: wmma.hpp:13