include/ck_tile/host/stream_utils.hpp Source File

• include
• ck_tile
• host

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:22

integer.hpp

ck_tile

Definition: cluster_descriptor.hpp:13

ck_tile::index_t

int32_t index_t

Definition: integer.hpp:9

hip_check_error.hpp