/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/3643/include/ck_tile/core/utility/persistent_async_input_scheduler.hpp Source File

• 3643
• include
• ck_tile
• core
• utility

persistent_async_input_scheduler.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 #include <cstdint>

7 

8 namespace ck_tile {

9 

26 struct PersistentAsyncInputScheduler

27 {

31  uint32_t tiles_per_chunk_m = 0;

32 

36  uint32_t* chunk_signals = nullptr;

37 

41  int32_t tile_idx_pivot_m = 0;

42 

46  uint32_t num_chunks = 0;

47 };

48 

49 } // namespace ck_tile

ck_tile

Definition: cluster_descriptor.hpp:13

ck_tile::int32_t

int32_t int32_t

Definition: integer.hpp:10

uint32_t

unsigned int uint32_t

Definition: stdint.h:126

ck_tile::PersistentAsyncInputScheduler

Scheduler for persistent GEMM kernels with asynchronous input streaming.

Definition: persistent_async_input_scheduler.hpp:27

ck_tile::PersistentAsyncInputScheduler::tiles_per_chunk_m

uint32_t tiles_per_chunk_m

Number of M-dimension tiles grouped into each chunk. Grouping tiles balances synchronization overhead...

Definition: persistent_async_input_scheduler.hpp:31

ck_tile::PersistentAsyncInputScheduler::tile_idx_pivot_m

int32_t tile_idx_pivot_m

Pivot offset for rotating the chunk assignment. Allows shifting which tiles map to which chunks,...

Definition: persistent_async_input_scheduler.hpp:41

ck_tile::PersistentAsyncInputScheduler::chunk_signals

uint32_t * chunk_signals

Device pointer to array of signal values (uint32_t), one per chunk. Producer sets signals to coordina...

Definition: persistent_async_input_scheduler.hpp:36

ck_tile::PersistentAsyncInputScheduler::num_chunks

uint32_t num_chunks

Number of signal chunks allocated. Must equal ceil((tiles_m + tile_idx_pivot_m) / tiles_per_chunk_m)....

Definition: persistent_async_input_scheduler.hpp:46