From 17be9aca75afcd2350b790fc63a78d0e02ebb034 Mon Sep 17 00:00:00 2001 From: scxfjiang Date: Thu, 9 Jan 2025 11:36:26 +0000 Subject: [PATCH] register gfx12xx --- xla/stream_executor/device_description.h | 12 +++++++++--- 1 file changed, 9 insertions(+), 3 deletions(-) diff --git a/xla/stream_executor/device_description.h b/xla/stream_executor/device_description.h index 195dd058aa64d..4d586a8b7e41e 100644 --- a/xla/stream_executor/device_description.h +++ b/xla/stream_executor/device_description.h @@ -202,6 +202,10 @@ class RocmComputeCapability { bool gfx11_rx7900() const { return gfx_version() == "gfx1100"; } + bool gfx1200() const { return gfx_version() == "gfx1200"; } + + bool gfx1201() const { return gfx_version() == "gfx1201"; } + bool has_nhwc_layout_support() const { return gfx9_mi100_or_later(); } bool has_bf16_dtype_support() const { return gfx9_mi100_or_later(); } @@ -227,9 +231,9 @@ class RocmComputeCapability { return gfx_version() != "gfx900" && gfx_version() != "gfx906"; } - bool has_hipblaslt() const { return gfx9_mi200_or_later(); } + bool has_hipblaslt() const { return gfx9_mi200_or_later() || gfx1200() || gfx1201(); } - bool has_fp8_support() const { return gfx9_mi300(); } + bool has_fp8_support() const { return gfx9_mi300() || gfx1200() || gfx1201(); } RocmComputeCapabilityProto ToProto() const { RocmComputeCapabilityProto proto; @@ -251,7 +255,9 @@ class RocmComputeCapability { "gfx90a", // MI200 "gfx940", "gfx941", "gfx942", // MI300 "gfx1030", // RX68xx / RX69xx - "gfx1100" // RX7900 + "gfx1100", // RX7900 + "gfx1200", + "gfx1201", }; };