Skip to content

feat(ck): enable gfx1170 and gfx1171 #13323

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

Closed
stsokolo wants to merge 1 commit into from
Closed

feat(ck): enable gfx1170 and gfx1171 #13323

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
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
4 changes: 2 additions & 2 deletions projects/composablekernel/CMakeLists.txt
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -403,9 +403,9 @@ if(NOT ENABLE_ASAN_PACKAGING)
# WORKAROUND: compiler does not yet fully support gfx12 targets, need to fix version above
set(CK_GPU_TARGETS "gfx908;gfx90a;gfx942;gfx1030;gfx1100;gfx1101;gfx1102;gfx1103")
elseif(NOT WIN32 AND ${hip_VERSION_FLAT} GREATER_EQUAL 600300000 AND ${hip_VERSION_FLAT} LESS 600400000)
set(CK_GPU_TARGETS "gfx908;gfx90a;gfx942;gfx1030;gfx1100;gfx1101;gfx1102;gfx1103;gfx1200;gfx1201")
set(CK_GPU_TARGETS "gfx908;gfx90a;gfx942;gfx1030;gfx1100;gfx1101;gfx1102;gfx1103;gfx1170;gfx1171;gfx1200;gfx1201")
elseif(NOT WIN32 AND ${hip_VERSION_FLAT} GREATER_EQUAL 600400000 AND ${hip_VERSION_FLAT} LESS 600443483)
set(CK_GPU_TARGETS "gfx908;gfx90a;gfx942;gfx1030;gfx1100;gfx1101;gfx1102;gfx1103;gfx1200;gfx1201;gfx950")
set(CK_GPU_TARGETS "gfx908;gfx90a;gfx942;gfx1030;gfx1100;gfx1101;gfx1102;gfx1103;gfx1170;gfx1171;gfx1200;gfx1201;gfx950")
elseif(NOT WIN32 AND ${hip_VERSION_FLAT} GREATER_EQUAL 600443483 AND ${hip_VERSION_FLAT} LESS 700200000)
set(CK_GPU_TARGETS "gfx908;gfx90a;gfx942;gfx950;gfx10-3-generic;gfx11-generic;gfx12-generic")
elseif(NOT WIN32 AND ${hip_VERSION_FLAT} GREATER_EQUAL 700200000)
Expand Down
3 changes: 2 additions & 1 deletion projects/composablekernel/include/ck/ck.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -88,7 +88,8 @@
#define CK_BUFFER_RESOURCE_3RD_DWORD 0x00020000
#elif defined(__gfx101__) || defined(__gfx103__)
#define CK_BUFFER_RESOURCE_3RD_DWORD 0x31014000
#elif defined(__gfx11__) || defined(__gfx120__)
#elif defined(__gfx11__) || defined(__gfx120__) || defined(__gfx1170__) || \
defined(__gfx1171__)
#define CK_BUFFER_RESOURCE_3RD_DWORD 0x31004000
#elif defined(__gfx125__)
#define CK_BUFFER_RESOURCE_3RD_DWORD 0
Expand Down
10 changes: 8 additions & 2 deletions projects/composablekernel/include/ck/host_utility/device_prop.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -71,6 +71,11 @@ inline int get_device_revision()
return props.asicRevision;
}

inline bool is_gfx117_supported()
{
return ck::get_device_name() == "gfx1170" || ck::get_device_name() == "gfx1171";
}

inline bool is_gfx12_supported()
{
return ck::get_device_name() == "gfx1200" || ck::get_device_name() == "gfx1201" ||
Expand Down Expand Up @@ -126,7 +131,7 @@ inline bool is_xdl_wmma_supported()
{
return true;
}
else if(is_gfx120_supported() || is_gfx11_supported())
else if(is_gfx120_supported() || is_gfx117_supported() || is_gfx11_supported())
{
if constexpr((MPerXDL32 != 16) || (NPerXDL32 != 16))
{
Expand Down Expand Up @@ -226,7 +231,8 @@ inline bool is_bf16_atomic_supported()

inline bool is_wmma_supported()
{
return is_gfx103_supported() || is_gfx11_supported() || is_gfx12_supported();
return is_gfx103_supported() || is_gfx11_supported() || is_gfx117_supported() ||
is_gfx12_supported();
}

inline bool is_tf32_supported()
Expand Down
3 changes: 2 additions & 1 deletion projects/composablekernel/include/ck_tile/core/config.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -292,7 +292,8 @@
#define CK_TILE_BUFFER_RESOURCE_3RD_DWORD 0x00020000
#elif defined(__gfx101__) || defined(__gfx103__) // for GPU code
#define CK_TILE_BUFFER_RESOURCE_3RD_DWORD 0x31014000
#elif defined(__gfx11__) || defined(__gfx120__)
#elif defined(__gfx11__) || defined(__gfx120__) || defined(__gfx1170__) || \
defined(__gfx1171__)
#define CK_TILE_BUFFER_RESOURCE_3RD_DWORD 0x31004000
#elif defined(__gfx125__)
#define CK_TILE_BUFFER_RESOURCE_3RD_DWORD 0x0
Expand Down
Loading