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

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

Composable Kernel: /home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/develop/include/ck_tile/host/device_prop.hpp Source File
device_prop.hpp
Go to the documentation of this file.
1 // SPDX-License-Identifier: MIT
2 // Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
3 
4 #pragma once
5 
6 #ifndef __HIPCC_RTC__
7 #include <string>
8 #include <string_view>
9 #include <hip/hip_runtime.h>
10 
11 namespace ck_tile {
12 
13 constexpr unsigned int fnv1a_hash(std::string_view str, unsigned int h = 2166136261u)
14 {
15  return str.empty() ? h
16  : fnv1a_hash(str.substr(1),
17  (h ^ static_cast<unsigned char>(str.front())) * 16777619u);
18 }
19 inline std::string get_device_name()
20 {
21  hipDeviceProp_t props{};
22  int device;
23  auto status = hipGetDevice(&device);
24  if(status != hipSuccess)
25  {
26  return std::string();
27  }
28  status = hipGetDeviceProperties(&props, device);
29  if(status != hipSuccess)
30  {
31  return std::string();
32  }
33  const std::string raw_name(props.gcnArchName);
34  const auto name = raw_name.substr(0, raw_name.find(':')); // str.substr(0, npos) returns str.
35  switch(fnv1a_hash(name))
36  {
37  // https://github.com/ROCm/MIOpen/blob/8498875aef84878e04c1eabefdf6571514891086/src/target_properties.cpp#L40
38  case fnv1a_hash("Ellesmere"):
39  case fnv1a_hash("Baffin"):
40  case fnv1a_hash("RacerX"):
41  case fnv1a_hash("Polaris10"):
42  case fnv1a_hash("Polaris11"):
43  case fnv1a_hash("Tonga"):
44  case fnv1a_hash("Fiji"):
45  case fnv1a_hash("gfx800"):
46  case fnv1a_hash("gfx802"):
47  case fnv1a_hash("gfx804"): return "gfx803";
48  case fnv1a_hash("Vega10"):
49  case fnv1a_hash("gfx901"): return "gfx900";
50  case fnv1a_hash("10.3.0 Sienna_Cichlid 18"): return "gfx1030";
51  default: return name;
52  }
53 }
54 
55 inline bool is_gfx11_supported()
56 {
57  return get_device_name() == "gfx1100" || get_device_name() == "gfx1101" ||
58  get_device_name() == "gfx1102" || get_device_name() == "gfx1103" ||
59  get_device_name() == "gfx1150" || get_device_name() == "gfx1151" ||
60  get_device_name() == "gfx1152";
61 }
62 
63 inline bool is_gfx12_supported()
64 {
65  return get_device_name() == "gfx1200" || get_device_name() == "gfx1201";
66 }
67 
68 inline bool is_load_tr_supported()
69 {
70  // Check if load transpose is supported.
71  return get_device_name() == "gfx950";
72 }
73 } // namespace ck_tile
74 
75 #endif
ck_tile
Definition: cluster_descriptor.hpp:13
ck_tile::fnv1a_hash
constexpr unsigned int fnv1a_hash(std::string_view str, unsigned int h=2166136261u)
Definition: device_prop.hpp:13
ck_tile::is_gfx12_supported
bool is_gfx12_supported()
Definition: device_prop.hpp:63
ck_tile::get_device_name
std::string get_device_name()
Definition: device_prop.hpp:19
ck_tile::is_load_tr_supported
bool is_load_tr_supported()
Definition: device_prop.hpp:68
ck_tile::is_gfx11_supported
bool is_gfx11_supported()
Definition: device_prop.hpp:55