/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/develop/include/ck_tile/ops/image_to_column/pipeline/tile_image_to_column_shape.hpp Source File

/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/develop/include/ck_tile/ops/image_to_column/pipeline/tile_image_to_column_shape.hpp Source File#

Composable Kernel: /home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/develop/include/ck_tile/ops/image_to_column/pipeline/tile_image_to_column_shape.hpp Source File

Go to the documentation of this file.

 // SPDX-License-Identifier: MIT
 // Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved.
  
 #pragma once
  
 #include "ck_tile/core.hpp"
  
 namespace ck_tile {
 template <typename ThreadTile, // Sequence<...
           typename WarpTile,   // Sequence<...
           typename BlockTile>  // Sequence<...
 struct TileImageToColumnShape
 {
     static constexpr index_t kMPerThread = ThreadTile::at(number<0>{});
     static constexpr index_t kKPerThread = ThreadTile::at(number<1>{});
  
     static constexpr index_t kMPerWarp       = WarpTile::at(number<0>{});
     static constexpr index_t kMThreadPerWarp = kMPerWarp / kMPerThread;
     static constexpr index_t kKThreadPerWarp = get_warp_size() / kMThreadPerWarp;
     static constexpr index_t kKPerWarp       = kKPerThread * kKThreadPerWarp;
  
     static constexpr index_t kMPerBlock = BlockTile::at(number<0>{});
     static constexpr index_t kKPerBlock = BlockTile::at(number<1>{});
  
     static constexpr index_t kMWarpPerBlock = kMPerBlock / kMPerWarp;
     static constexpr index_t kKWarpPerBlock = kKPerBlock / kKPerWarp;
  
     static constexpr index_t kBlockSize = get_warp_size() * kMWarpPerBlock * kKWarpPerBlock;
 };
  
 } // namespace ck_tile