Skip to content

[SYCL][NFC] Refactor device architecture data structures #12113

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 2 commits into from
Dec 8, 2023
Merged

[SYCL][NFC] Refactor device architecture data structures #12113

Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
9c2ab0d
[SYCL][NFC] Refactor device architecture data structures
dm-vodopyanov Dec 7, 2023
4c1d549
Remove extra code
dm-vodopyanov Dec 7, 2023
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
164 changes: 84 additions & 80 deletions sycl/source/detail/device_info.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -557,94 +557,92 @@ struct get_device_info_impl<range<Dimensions>,
}
};

// This macro is only for AMD and NVIDIA GPU architectures
#define NVIDIA_AMD_ARCHES(X) \
X("5.0", oneapi_exp_arch::nvidia_gpu_sm_50) \
X("5.2", oneapi_exp_arch::nvidia_gpu_sm_52) \
X("5.3", oneapi_exp_arch::nvidia_gpu_sm_53) \
X("6.0", oneapi_exp_arch::nvidia_gpu_sm_60) \
X("6.1", oneapi_exp_arch::nvidia_gpu_sm_61) \
X("6.2", oneapi_exp_arch::nvidia_gpu_sm_62) \
X("7.0", oneapi_exp_arch::nvidia_gpu_sm_70) \
X("7.2", oneapi_exp_arch::nvidia_gpu_sm_72) \
X("7.5", oneapi_exp_arch::nvidia_gpu_sm_75) \
X("8.0", oneapi_exp_arch::nvidia_gpu_sm_80) \
X("8.6", oneapi_exp_arch::nvidia_gpu_sm_86) \
X("8.7", oneapi_exp_arch::nvidia_gpu_sm_87) \
X("8.9", oneapi_exp_arch::nvidia_gpu_sm_89) \
X("9.0", oneapi_exp_arch::nvidia_gpu_sm_90) \
X("gfx701", oneapi_exp_arch::amd_gpu_gfx701) \
X("gfx702", oneapi_exp_arch::amd_gpu_gfx702) \
X("gfx801", oneapi_exp_arch::amd_gpu_gfx801) \
X("gfx802", oneapi_exp_arch::amd_gpu_gfx802) \
X("gfx803", oneapi_exp_arch::amd_gpu_gfx803) \
X("gfx805", oneapi_exp_arch::amd_gpu_gfx805) \
X("gfx810", oneapi_exp_arch::amd_gpu_gfx810) \
X("gfx900", oneapi_exp_arch::amd_gpu_gfx900) \
X("gfx902", oneapi_exp_arch::amd_gpu_gfx902) \
X("gfx904", oneapi_exp_arch::amd_gpu_gfx904) \
X("gfx906", oneapi_exp_arch::amd_gpu_gfx906) \
X("gfx908", oneapi_exp_arch::amd_gpu_gfx908) \
X("gfx90a", oneapi_exp_arch::amd_gpu_gfx90a) \
X("gfx1010", oneapi_exp_arch::amd_gpu_gfx1010) \
X("gfx1011", oneapi_exp_arch::amd_gpu_gfx1011) \
X("gfx1012", oneapi_exp_arch::amd_gpu_gfx1012) \
X("gfx1013", oneapi_exp_arch::amd_gpu_gfx1013) \
X("gfx1030", oneapi_exp_arch::amd_gpu_gfx1030) \
X("gfx1031", oneapi_exp_arch::amd_gpu_gfx1031) \
X("gfx1032", oneapi_exp_arch::amd_gpu_gfx1032) \
X("gfx1034", oneapi_exp_arch::amd_gpu_gfx1034)

// This macro is only for Intel GPU architectures
#define INTEL_GPU_ARCHES(X) \
X(0x02000000, oneapi_exp_arch::intel_gpu_bdw) \
X(0x02400009, oneapi_exp_arch::intel_gpu_skl) \
X(0x02404009, oneapi_exp_arch::intel_gpu_kbl) \
X(0x02408009, oneapi_exp_arch::intel_gpu_cfl) \
X(0x0240c000, oneapi_exp_arch::intel_gpu_apl) \
X(0x02410000, oneapi_exp_arch::intel_gpu_glk) \
X(0x02414000, oneapi_exp_arch::intel_gpu_whl) \
X(0x02418000, oneapi_exp_arch::intel_gpu_aml) \
X(0x0241c000, oneapi_exp_arch::intel_gpu_cml) \
X(0x02c00000, oneapi_exp_arch::intel_gpu_icllp) \
X(0x02c08000, oneapi_exp_arch::intel_gpu_ehl) \
X(0x03000000, oneapi_exp_arch::intel_gpu_tgllp) \
X(0x03004000, oneapi_exp_arch::intel_gpu_rkl) \
X(0x03008000, oneapi_exp_arch::intel_gpu_adl_s) \
X(0x0300c000, oneapi_exp_arch::intel_gpu_adl_p) \
X(0x03010000, oneapi_exp_arch::intel_gpu_adl_n) \
X(0x03028000, oneapi_exp_arch::intel_gpu_dg1) \
X(0x030dc008, oneapi_exp_arch::intel_gpu_acm_g10) \
X(0x030e0005, oneapi_exp_arch::intel_gpu_acm_g11) \
X(0x030e4000, oneapi_exp_arch::intel_gpu_acm_g12) \
X(0x030f0007, oneapi_exp_arch::intel_gpu_pvc)

// This macro is only for Intel CPU architectures
// TODO: extend the macro with other CPU architectures when they will be added
// to ext_oneapi_device_architecture
#define INTEL_CPU_ARCHES(X) \
X(8, oneapi_exp_arch::intel_cpu_spr) \
X(9, oneapi_exp_arch::intel_cpu_gnr)

#define CMP_NVIDIA_AMD(s, i) \
if (strcmp(s, arch) == 0) \
return i;

#define CMP_INTEL(p, i) \
if (p == arch) \
return i;
using oneapi_exp_arch = sycl::ext::oneapi::experimental::architecture;

// Only for NVIDIA and AMD GPU architectures
constexpr std::pair<const char *, oneapi_exp_arch> NvidiaAmdGPUArchitectures[] =
{
{"5.0", oneapi_exp_arch::nvidia_gpu_sm_50},
{"5.2", oneapi_exp_arch::nvidia_gpu_sm_52},
{"5.3", oneapi_exp_arch::nvidia_gpu_sm_53},
{"6.0", oneapi_exp_arch::nvidia_gpu_sm_60},
{"6.1", oneapi_exp_arch::nvidia_gpu_sm_61},
{"6.2", oneapi_exp_arch::nvidia_gpu_sm_62},
{"7.0", oneapi_exp_arch::nvidia_gpu_sm_70},
{"7.2", oneapi_exp_arch::nvidia_gpu_sm_72},
{"7.5", oneapi_exp_arch::nvidia_gpu_sm_75},
{"8.0", oneapi_exp_arch::nvidia_gpu_sm_80},
{"8.6", oneapi_exp_arch::nvidia_gpu_sm_86},
{"8.7", oneapi_exp_arch::nvidia_gpu_sm_87},
{"8.9", oneapi_exp_arch::nvidia_gpu_sm_89},
{"9.0", oneapi_exp_arch::nvidia_gpu_sm_90},
{"gfx701", oneapi_exp_arch::amd_gpu_gfx701},
{"gfx702", oneapi_exp_arch::amd_gpu_gfx702},
{"gfx801", oneapi_exp_arch::amd_gpu_gfx801},
{"gfx802", oneapi_exp_arch::amd_gpu_gfx802},
{"gfx803", oneapi_exp_arch::amd_gpu_gfx803},
{"gfx805", oneapi_exp_arch::amd_gpu_gfx805},
{"gfx810", oneapi_exp_arch::amd_gpu_gfx810},
{"gfx900", oneapi_exp_arch::amd_gpu_gfx900},
{"gfx902", oneapi_exp_arch::amd_gpu_gfx902},
{"gfx904", oneapi_exp_arch::amd_gpu_gfx904},
{"gfx906", oneapi_exp_arch::amd_gpu_gfx906},
{"gfx908", oneapi_exp_arch::amd_gpu_gfx908},
{"gfx90a", oneapi_exp_arch::amd_gpu_gfx90a},
{"gfx1010", oneapi_exp_arch::amd_gpu_gfx1010},
{"gfx1011", oneapi_exp_arch::amd_gpu_gfx1011},
{"gfx1012", oneapi_exp_arch::amd_gpu_gfx1012},
{"gfx1013", oneapi_exp_arch::amd_gpu_gfx1013},
{"gfx1030", oneapi_exp_arch::amd_gpu_gfx1030},
{"gfx1031", oneapi_exp_arch::amd_gpu_gfx1031},
{"gfx1032", oneapi_exp_arch::amd_gpu_gfx1032},
{"gfx1034", oneapi_exp_arch::amd_gpu_gfx1034},
};

// Only for Intel GPU architectures
constexpr std::pair<const int, oneapi_exp_arch> IntelGPUArchitectures[] = {
{0x02000000, oneapi_exp_arch::intel_gpu_bdw},
{0x02400009, oneapi_exp_arch::intel_gpu_skl},
{0x02404009, oneapi_exp_arch::intel_gpu_kbl},
{0x02408009, oneapi_exp_arch::intel_gpu_cfl},
{0x0240c000, oneapi_exp_arch::intel_gpu_apl},
{0x02410000, oneapi_exp_arch::intel_gpu_glk},
{0x02414000, oneapi_exp_arch::intel_gpu_whl},
{0x02418000, oneapi_exp_arch::intel_gpu_aml},
{0x0241c000, oneapi_exp_arch::intel_gpu_cml},
{0x02c00000, oneapi_exp_arch::intel_gpu_icllp},
{0x02c08000, oneapi_exp_arch::intel_gpu_ehl},
{0x03000000, oneapi_exp_arch::intel_gpu_tgllp},
{0x03004000, oneapi_exp_arch::intel_gpu_rkl},
{0x03008000, oneapi_exp_arch::intel_gpu_adl_s},
{0x0300c000, oneapi_exp_arch::intel_gpu_adl_p},
{0x03010000, oneapi_exp_arch::intel_gpu_adl_n},
{0x03028000, oneapi_exp_arch::intel_gpu_dg1},
{0x030dc008, oneapi_exp_arch::intel_gpu_acm_g10},
{0x030e0005, oneapi_exp_arch::intel_gpu_acm_g11},
{0x030e4000, oneapi_exp_arch::intel_gpu_acm_g12},
{0x030f0007, oneapi_exp_arch::intel_gpu_pvc},
};

// Only for Intel CPU architectures
constexpr std::pair<const int, oneapi_exp_arch> IntelCPUArchitectures[] = {
{8, oneapi_exp_arch::intel_cpu_spr},
{9, oneapi_exp_arch::intel_cpu_gnr},
};

template <>
struct get_device_info_impl<
ext::oneapi::experimental::architecture,
ext::oneapi::experimental::info::device::architecture> {
static ext::oneapi::experimental::architecture get(const DeviceImplPtr &Dev) {
using oneapi_exp_arch = sycl::ext::oneapi::experimental::architecture;
backend CurrentBackend = Dev->getBackend();
if (Dev->is_gpu() && (backend::ext_oneapi_level_zero == CurrentBackend ||
backend::opencl == CurrentBackend)) {
auto MapArchIDToArchName = [](const int arch) {
INTEL_GPU_ARCHES(CMP_INTEL);
for (const auto &Item : IntelGPUArchitectures) {
if (Item.first == arch)
return Item.second;
}
throw sycl::exception(
make_error_code(errc::runtime),
"The current device architecture is not supported by "
Expand All @@ -660,7 +658,10 @@ struct get_device_info_impl<
} else if (Dev->is_gpu() && (backend::ext_oneapi_cuda == CurrentBackend ||
backend::ext_oneapi_hip == CurrentBackend)) {
auto MapArchIDToArchName = [](const char *arch) {
NVIDIA_AMD_ARCHES(CMP_NVIDIA_AMD);
for (const auto &Item : NvidiaAmdGPUArchitectures) {
if (std::string_view(Item.first) == arch)
return Item.second;
}
throw sycl::exception(
make_error_code(errc::runtime),
"The current device architecture is not supported by "
Expand All @@ -680,7 +681,10 @@ struct get_device_info_impl<
return MapArchIDToArchName(DeviceArchSubstr.data());
} else if (Dev->is_cpu() && backend::opencl == CurrentBackend) {
auto MapArchIDToArchName = [](const int arch) {
INTEL_CPU_ARCHES(CMP_INTEL);
for (const auto &Item : IntelCPUArchitectures) {
if (Item.first == arch)
return Item.second;
}
return sycl::ext::oneapi::experimental::architecture::x86_64;
};
uint32_t DeviceIp;
Expand Down