/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/develop/include/ck_tile/host/stream_utils.hpp Source File

/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/develop/include/ck_tile/host/stream_utils.hpp Source File#

Composable Kernel: /home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/develop/include/ck_tile/host/stream_utils.hpp Source File
stream_utils.hpp
Go to the documentation of this file.
1 // SPDX-License-Identifier: MIT
2 // Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved.
3 
4 #pragma once
5 
6 #include <hip/hip_runtime_api.h>
7 
8 #include "ck_tile/core/numeric/integer.hpp"
9 #include "ck_tile/host/stream_config.hpp"
10 #include "ck_tile/host/hip_check_error.hpp"
11 
12 namespace ck_tile {
13 
14 static inline index_t get_available_compute_units(const stream_config& s)
15 {
16  constexpr static uint32_t MAX_MASK_DWORDS = 64;
17 
18  // assume at most 64*32 = 2048 CUs
19  uint32_t cu_mask[MAX_MASK_DWORDS]{};
20 
21  auto count_set_bits = [](uint32_t dword) {
22  index_t count = 0;
23  while(dword != 0)
24  {
25  if(dword & 0x1)
26  {
27  count++;
28  }
29  dword = dword >> 1;
30  }
31  return count;
32  };
33 
34  HIP_CHECK_ERROR(hipExtStreamGetCUMask(s.stream_id_, MAX_MASK_DWORDS, &cu_mask[0]));
35 
36  index_t num_cu = 0;
37  for(uint32_t i = 0; i < MAX_MASK_DWORDS; i++)
38  {
39  num_cu += count_set_bits(cu_mask[i]);
40  }
41 
42  return num_cu;
43 };
44 
45 } // namespace ck_tile
stream_config.hpp
HIP_CHECK_ERROR
#define HIP_CHECK_ERROR(retval_or_funcall)
Definition: hip_check_error.hpp:21
ck_tile
Definition: cluster_descriptor.hpp:13
ck_tile::index_t
int32_t index_t
Definition: integer.hpp:9
uint32_t
unsigned int uint32_t
Definition: stdint.h:126
hip_check_error.hpp