Skip to content

Commit 17d3a5e

Browse files
harkgill-amdillsilin
authored andcommitted
Add missing gfx1033 to gfx103 group definition in ck (#5141)
## Motivation Resolving PyTorch build failures when enabling builds for gfx103X-all family in TheRock. ROCm/TheRock#3763. `gfx1033` is the only failing architecture in the family and the failures point to missing support in CK. ## Technical Details PyTorch build fails with repeated error message ``` /__w/TheRock/TheRock/external-builds/pytorch/pytorch/aten/src/ATen/../../../third_party/composable_kernel/include/ck/utility/amd_buffer_addressing_builtins.hpp:33:48: error: use of undeclared identifier 'CK_BUFFER_RESOURCE_3RD_DWORD' 33 | wave_buffer_resource.config(Number<3>{}) = CK_BUFFER_RESOURCE_3RD_DWORD; | ^~~~~~~~~~~~~~~~~~~~~~~~~~~~ ``` `gfx1033` is missing from the `__gfx103__` group which results in `CK_BUFFER_RESOURCE_3RD_DWORD` never being defined for it. Adding in `gfx1033` to the missing files which should be the minimum fix to allow torch builds to pass. ## Test Plan Compile sample test file and target gfx1033 ``` ... #ifdef __HIP_DEVICE_COMPILE__ static_assert(CK_BUFFER_RESOURCE_3RD_DWORD == 0x31014000, "wrong device value"); #else static_assert(CK_BUFFER_RESOURCE_3RD_DWORD == -1, "wrong host value"); #endif ``` ## Test Result Prior to the applying patch, compilation fails with `error: use of undeclared identifier 'CK_BUFFER_RESOURCE_3RD_DWORD'` After applying patch, test file compiles successfully. ## Submission Checklist - [X] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests. --------- Co-authored-by: Illia Silin <[email protected]>
1 parent d8b7a2a commit 17d3a5e

File tree

4 files changed

+19
-6
lines changed

4 files changed

+19
-6
lines changed

projects/composablekernel/include/ck/ck.hpp

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -63,8 +63,8 @@
6363
#define __gfx101__
6464
#endif
6565
#if defined(__gfx1030__) || defined(__gfx1031__) || defined(__gfx1032__) || \
66-
defined(__gfx1034__) || defined(__gfx1035__) || defined(__gfx1036__) || \
67-
defined(__gfx10_3_generic__)
66+
defined(__gfx1033__) || defined(__gfx1034__) || defined(__gfx1035__) || \
67+
defined(__gfx1036__) || defined(__gfx10_3_generic__)
6868
#define __gfx103__
6969
#endif
7070
#if defined(__gfx1100__) || defined(__gfx1101__) || defined(__gfx1102__) || \

projects/composablekernel/include/ck/host_utility/device_prop.hpp

Lines changed: 3 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -125,8 +125,9 @@ inline bool is_gfx101_supported()
125125
inline bool is_gfx103_supported()
126126
{
127127
return ck::get_device_name() == "gfx1030" || ck::get_device_name() == "gfx1031" ||
128-
ck::get_device_name() == "gfx1032" || ck::get_device_name() == "gfx1034" ||
129-
ck::get_device_name() == "gfx1035" || ck::get_device_name() == "gfx1036";
128+
ck::get_device_name() == "gfx1032" || ck::get_device_name() == "gfx1033" ||
129+
ck::get_device_name() == "gfx1034" || ck::get_device_name() == "gfx1035" ||
130+
ck::get_device_name() == "gfx1036";
130131
}
131132

132133
inline bool is_wmma_supported()

projects/composablekernel/include/ck_tile/core/arch/arch.hpp

Lines changed: 5 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -88,6 +88,7 @@ enum struct amdgcn_target_id
8888
GFX1030 = 0x1030,
8989
GFX1031 = 0x1031,
9090
GFX1032 = 0x1032,
91+
GFX1033 = 0x1033,
9192
GFX1034 = 0x1034,
9293
GFX1035 = 0x1035,
9394
GFX1036 = 0x1036,
@@ -284,6 +285,7 @@ constexpr auto get_compiler_target()
284285
MAP_COMPILER_STATE_TO_GFX10_3_TARGET(CK_TILE_ARCH_GFX1030, GFX1030);
285286
MAP_COMPILER_STATE_TO_GFX10_3_TARGET(CK_TILE_ARCH_GFX1031, GFX1031);
286287
MAP_COMPILER_STATE_TO_GFX10_3_TARGET(CK_TILE_ARCH_GFX1032, GFX1032);
288+
MAP_COMPILER_STATE_TO_GFX10_3_TARGET(CK_TILE_ARCH_GFX1033, GFX1033);
287289
MAP_COMPILER_STATE_TO_GFX10_3_TARGET(CK_TILE_ARCH_GFX1034, GFX1034);
288290
MAP_COMPILER_STATE_TO_GFX10_3_TARGET(CK_TILE_ARCH_GFX1035, GFX1035);
289291
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*
351353
MAP_HIP_DEVICE_PROP_GCN_ARCH_NAME_STRING_TO_TARGET_ID("gfx1030", GFX1030);
352354
MAP_HIP_DEVICE_PROP_GCN_ARCH_NAME_STRING_TO_TARGET_ID("gfx1031", GFX1031);
353355
MAP_HIP_DEVICE_PROP_GCN_ARCH_NAME_STRING_TO_TARGET_ID("gfx1032", GFX1032);
356+
MAP_HIP_DEVICE_PROP_GCN_ARCH_NAME_STRING_TO_TARGET_ID("gfx1033", GFX1033);
354357
MAP_HIP_DEVICE_PROP_GCN_ARCH_NAME_STRING_TO_TARGET_ID("gfx1034", GFX1034);
355358
MAP_HIP_DEVICE_PROP_GCN_ARCH_NAME_STRING_TO_TARGET_ID("gfx1035", GFX1035);
356359
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()
607610
MAP_COMPILER_STATE_TO_GFX10_3_TARGET(CK_TILE_ARCH_GFX1030, GFX1030);
608611
MAP_COMPILER_STATE_TO_GFX10_3_TARGET(CK_TILE_ARCH_GFX1031, GFX1031);
609612
MAP_COMPILER_STATE_TO_GFX10_3_TARGET(CK_TILE_ARCH_GFX1032, GFX1032);
613+
MAP_COMPILER_STATE_TO_GFX10_3_TARGET(CK_TILE_ARCH_GFX1033, GFX1033);
610614
MAP_COMPILER_STATE_TO_GFX10_3_TARGET(CK_TILE_ARCH_GFX1034, GFX1034);
611615
MAP_COMPILER_STATE_TO_GFX10_3_TARGET(CK_TILE_ARCH_GFX1035, GFX1035);
612616
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
688692
MAP_HIP_DEVICE_PROP_GCN_ARCH_NAME_STRING_TO_GFX10_3_TARGET("gfx1030", GFX1030);
689693
MAP_HIP_DEVICE_PROP_GCN_ARCH_NAME_STRING_TO_GFX10_3_TARGET("gfx1031", GFX1031);
690694
MAP_HIP_DEVICE_PROP_GCN_ARCH_NAME_STRING_TO_GFX10_3_TARGET("gfx1032", GFX1032);
695+
MAP_HIP_DEVICE_PROP_GCN_ARCH_NAME_STRING_TO_GFX10_3_TARGET("gfx1033", GFX1033);
691696
MAP_HIP_DEVICE_PROP_GCN_ARCH_NAME_STRING_TO_GFX10_3_TARGET("gfx1034", GFX1034);
692697
MAP_HIP_DEVICE_PROP_GCN_ARCH_NAME_STRING_TO_GFX10_3_TARGET("gfx1035", GFX1035);
693698
MAP_HIP_DEVICE_PROP_GCN_ARCH_NAME_STRING_TO_GFX10_3_TARGET("gfx1036", GFX1036);

projects/composablekernel/include/ck_tile/core/config.hpp

Lines changed: 9 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -15,8 +15,8 @@
1515
#define __gfx101__
1616
#endif
1717
#if defined(__gfx1030__) || defined(__gfx1031__) || defined(__gfx1032__) || \
18-
defined(__gfx1034__) || defined(__gfx1035__) || defined(__gfx1036__) || \
19-
defined(__gfx10_3_generic__)
18+
defined(__gfx1033__) || defined(__gfx1034__) || defined(__gfx1035__) || \
19+
defined(__gfx1036__) || defined(__gfx10_3_generic__)
2020
#define __gfx103__
2121
#endif
2222
#if defined(__gfx1100__) || defined(__gfx1101__) || defined(__gfx1102__) || \
@@ -405,6 +405,12 @@ struct amdgcn_compiler_target_state
405405
static constexpr bool CK_TILE_ARCH_GFX1032 = false;
406406
#endif // __gfx1032__
407407

408+
#if defined(__gfx1033__)
409+
static constexpr bool CK_TILE_ARCH_GFX1033 = true;
410+
#else
411+
static constexpr bool CK_TILE_ARCH_GFX1033 = false;
412+
#endif // __gfx1033__
413+
408414
#if defined(__gfx1034__)
409415
static constexpr bool CK_TILE_ARCH_GFX1034 = true;
410416
#else
@@ -537,6 +543,7 @@ CK_TILE_HOST_DEVICE static constexpr uint32_t count_values_of(T search, Ts... se
537543
amdgcn_compiler_target_state::CK_TILE_ARCH_GFX1030, \
538544
amdgcn_compiler_target_state::CK_TILE_ARCH_GFX1031, \
539545
amdgcn_compiler_target_state::CK_TILE_ARCH_GFX1032, \
546+
amdgcn_compiler_target_state::CK_TILE_ARCH_GFX1033, \
540547
amdgcn_compiler_target_state::CK_TILE_ARCH_GFX1034, \
541548
amdgcn_compiler_target_state::CK_TILE_ARCH_GFX1035, \
542549
amdgcn_compiler_target_state::CK_TILE_ARCH_GFX1036, \

0 commit comments

Comments
 (0)