diff --git a/include/ck/ck.hpp b/include/ck/ck.hpp index 879fb31ca5..7448c7a31a 100644 --- a/include/ck/ck.hpp +++ b/include/ck/ck.hpp @@ -63,8 +63,8 @@ #define __gfx101__ #endif #if defined(__gfx1030__) || defined(__gfx1031__) || defined(__gfx1032__) || \ - defined(__gfx1034__) || defined(__gfx1035__) || defined(__gfx1036__) || \ - defined(__gfx10_3_generic__) + defined(__gfx1033__) || defined(__gfx1034__) || defined(__gfx1035__) || \ + defined(__gfx1036__) || defined(__gfx10_3_generic__) #define __gfx103__ #endif #if defined(__gfx1100__) || defined(__gfx1101__) || defined(__gfx1102__) || \ diff --git a/include/ck/host_utility/device_prop.hpp b/include/ck/host_utility/device_prop.hpp index 7191ad2c8a..97852531a9 100644 --- a/include/ck/host_utility/device_prop.hpp +++ b/include/ck/host_utility/device_prop.hpp @@ -125,8 +125,9 @@ inline bool is_gfx101_supported() inline bool is_gfx103_supported() { return ck::get_device_name() == "gfx1030" || ck::get_device_name() == "gfx1031" || - ck::get_device_name() == "gfx1032" || ck::get_device_name() == "gfx1034" || - ck::get_device_name() == "gfx1035" || ck::get_device_name() == "gfx1036"; + ck::get_device_name() == "gfx1032" || ck::get_device_name() == "gfx1033" || + ck::get_device_name() == "gfx1034" || ck::get_device_name() == "gfx1035" || + ck::get_device_name() == "gfx1036"; } inline bool is_wmma_supported() diff --git a/include/ck_tile/core/arch/arch.hpp b/include/ck_tile/core/arch/arch.hpp index 62d7971a8a..0775b34eef 100644 --- a/include/ck_tile/core/arch/arch.hpp +++ b/include/ck_tile/core/arch/arch.hpp @@ -88,6 +88,7 @@ enum struct amdgcn_target_id GFX1030 = 0x1030, GFX1031 = 0x1031, GFX1032 = 0x1032, + GFX1033 = 0x1033, GFX1034 = 0x1034, GFX1035 = 0x1035, GFX1036 = 0x1036, @@ -284,6 +285,7 @@ constexpr auto get_compiler_target() MAP_COMPILER_STATE_TO_GFX10_3_TARGET(CK_TILE_ARCH_GFX1030, GFX1030); MAP_COMPILER_STATE_TO_GFX10_3_TARGET(CK_TILE_ARCH_GFX1031, GFX1031); MAP_COMPILER_STATE_TO_GFX10_3_TARGET(CK_TILE_ARCH_GFX1032, GFX1032); + MAP_COMPILER_STATE_TO_GFX10_3_TARGET(CK_TILE_ARCH_GFX1033, GFX1033); MAP_COMPILER_STATE_TO_GFX10_3_TARGET(CK_TILE_ARCH_GFX1034, GFX1034); MAP_COMPILER_STATE_TO_GFX10_3_TARGET(CK_TILE_ARCH_GFX1035, GFX1035); MAP_COMPILER_STATE_TO_GFX10_3_TARGET(CK_TILE_ARCH_GFX1036, GFX1036); @@ -351,6 +353,7 @@ CK_TILE_HOST auto hip_device_prop_gcn_arch_name_to_amdgcn_target_id(char const* MAP_HIP_DEVICE_PROP_GCN_ARCH_NAME_STRING_TO_TARGET_ID("gfx1030", GFX1030); MAP_HIP_DEVICE_PROP_GCN_ARCH_NAME_STRING_TO_TARGET_ID("gfx1031", GFX1031); MAP_HIP_DEVICE_PROP_GCN_ARCH_NAME_STRING_TO_TARGET_ID("gfx1032", GFX1032); + MAP_HIP_DEVICE_PROP_GCN_ARCH_NAME_STRING_TO_TARGET_ID("gfx1033", GFX1033); MAP_HIP_DEVICE_PROP_GCN_ARCH_NAME_STRING_TO_TARGET_ID("gfx1034", GFX1034); MAP_HIP_DEVICE_PROP_GCN_ARCH_NAME_STRING_TO_TARGET_ID("gfx1035", GFX1035); MAP_HIP_DEVICE_PROP_GCN_ARCH_NAME_STRING_TO_TARGET_ID("gfx1036", GFX1036); @@ -607,6 +610,7 @@ CK_TILE_HOST_DEVICE constexpr auto get_compiler_target() MAP_COMPILER_STATE_TO_GFX10_3_TARGET(CK_TILE_ARCH_GFX1030, GFX1030); MAP_COMPILER_STATE_TO_GFX10_3_TARGET(CK_TILE_ARCH_GFX1031, GFX1031); MAP_COMPILER_STATE_TO_GFX10_3_TARGET(CK_TILE_ARCH_GFX1032, GFX1032); + MAP_COMPILER_STATE_TO_GFX10_3_TARGET(CK_TILE_ARCH_GFX1033, GFX1033); MAP_COMPILER_STATE_TO_GFX10_3_TARGET(CK_TILE_ARCH_GFX1034, GFX1034); MAP_COMPILER_STATE_TO_GFX10_3_TARGET(CK_TILE_ARCH_GFX1035, GFX1035); MAP_COMPILER_STATE_TO_GFX10_3_TARGET(CK_TILE_ARCH_GFX1036, GFX1036); @@ -688,6 +692,7 @@ CK_TILE_HOST auto hip_device_prop_gcn_arch_name_to_amdgcn_target(char const* tes MAP_HIP_DEVICE_PROP_GCN_ARCH_NAME_STRING_TO_GFX10_3_TARGET("gfx1030", GFX1030); MAP_HIP_DEVICE_PROP_GCN_ARCH_NAME_STRING_TO_GFX10_3_TARGET("gfx1031", GFX1031); MAP_HIP_DEVICE_PROP_GCN_ARCH_NAME_STRING_TO_GFX10_3_TARGET("gfx1032", GFX1032); + MAP_HIP_DEVICE_PROP_GCN_ARCH_NAME_STRING_TO_GFX10_3_TARGET("gfx1033", GFX1033); MAP_HIP_DEVICE_PROP_GCN_ARCH_NAME_STRING_TO_GFX10_3_TARGET("gfx1034", GFX1034); MAP_HIP_DEVICE_PROP_GCN_ARCH_NAME_STRING_TO_GFX10_3_TARGET("gfx1035", GFX1035); MAP_HIP_DEVICE_PROP_GCN_ARCH_NAME_STRING_TO_GFX10_3_TARGET("gfx1036", GFX1036); diff --git a/include/ck_tile/core/config.hpp b/include/ck_tile/core/config.hpp index a057ae9052..036e241c95 100644 --- a/include/ck_tile/core/config.hpp +++ b/include/ck_tile/core/config.hpp @@ -15,8 +15,8 @@ #define __gfx101__ #endif #if defined(__gfx1030__) || defined(__gfx1031__) || defined(__gfx1032__) || \ - defined(__gfx1034__) || defined(__gfx1035__) || defined(__gfx1036__) || \ - defined(__gfx10_3_generic__) + defined(__gfx1033__) || defined(__gfx1034__) || defined(__gfx1035__) || \ + defined(__gfx1036__) || defined(__gfx10_3_generic__) #define __gfx103__ #endif #if defined(__gfx1100__) || defined(__gfx1101__) || defined(__gfx1102__) || \ @@ -405,6 +405,12 @@ struct amdgcn_compiler_target_state static constexpr bool CK_TILE_ARCH_GFX1032 = false; #endif // __gfx1032__ +#if defined(__gfx1033__) + static constexpr bool CK_TILE_ARCH_GFX1033 = true; +#else + static constexpr bool CK_TILE_ARCH_GFX1033 = false; +#endif // __gfx1033__ + #if defined(__gfx1034__) static constexpr bool CK_TILE_ARCH_GFX1034 = true; #else @@ -537,6 +543,7 @@ CK_TILE_HOST_DEVICE static constexpr uint32_t count_values_of(T search, Ts... se amdgcn_compiler_target_state::CK_TILE_ARCH_GFX1030, \ amdgcn_compiler_target_state::CK_TILE_ARCH_GFX1031, \ amdgcn_compiler_target_state::CK_TILE_ARCH_GFX1032, \ + amdgcn_compiler_target_state::CK_TILE_ARCH_GFX1033, \ amdgcn_compiler_target_state::CK_TILE_ARCH_GFX1034, \ amdgcn_compiler_target_state::CK_TILE_ARCH_GFX1035, \ amdgcn_compiler_target_state::CK_TILE_ARCH_GFX1036, \