Fix issue with multiple targets and remove smfmac tests from unsupported test targets #1372

junliume · 2024-07-03T01:53:23Z

junliume · 2024-07-03T01:54:40Z

@illsilin @jakpiase please help to check if this check if complete, build multiple targets and tests multiple targets for CK is indeed tricky

carlushuang

LGTM

junliume · 2024-07-03T03:10:20Z

include/ck/utility/amd_smfmac.hpp

        reg_c.template AsType<float4_t>()(Number<0>{}) = __builtin_amdgcn_smfmac_f32_16x16x32_f16(
            reg_a, reg_b, reg_c.template AsType<float4_t>()[Number<0>{}], reg_idx, 0, 0);
+#else
+        ignore = reg_a;


These changes are consolidated from #1358

junliume · 2024-07-03T03:40:16Z

test/grouped_convnd_bwd_data/CMakeLists.txt

@@ -2,11 +2,11 @@ add_gtest_executable(test_grouped_convnd_bwd_data test_grouped_convnd_bwd_data_x
 if(result EQUAL 0)
    target_link_libraries(test_grouped_convnd_bwd_data PRIVATE utility device_grouped_conv2d_bwd_data_instance device_grouped_conv3d_bwd_data_instance)
 endif()
-add_gtest_executable(test_grouped_convnd_bwd_data_interface test_grouped_convnd_bwd_data_interface_xdl.cpp)


For reviewers: these are fixing duplicate test name issue

include/ck/utility/amd_smfmac.hpp

jakpiase

From what I remember smfmac instructions won't work on earlier ROCm versions than 6.2. Is there any way to stop using this test on cmake level based on that condition?

junliume · 2024-07-03T08:04:39Z

From what I remember smfmac instructions won't work on earlier ROCm versions than 6.2. Is there any way to stop using this test on cmake level based on that condition?

maybe if(CK_hip_VERSION VERSION_GREATER_EQUAL 6.2) ?

junliume · 2024-07-03T08:06:45Z

@jakpiase @aosewski please feel free to edit this branch in this PR :) I will be offline for the next a few hours.

All the remaining problems are easily reproducible by:

CXX=/opt/rocm/bin/amdclang++ cmake -DCMAKE_PREFIX_PATH=/opt/rocm -DCMAKE_BUILD_TYPE=Release -DGPU_TARGETS="gfx1100;gfx90a;gfx942" ..

and then

make -j$(nproc) check

… time

client_example/25_wrapper/CMakeLists.txt

atamazov · 2024-07-03T22:04:51Z

client_example/25_wrapper/wrapper_basic_gemm.cpp

+#if(!defined(__HIP_DEVICE_COMPILE__) || defined(__gfx908__) || defined(__gfx90a__) || \
+    defined(__gfx94__))


@junliume The expected effect is to compile the code for the host and for MI100/200/300, right? Looks good (at least for clang 19, see https://clang.llvm.org/docs/AMDGPUSupport.html) except __gfx94__ which looks suspicious. If you are in doubt too, then you can try this:

#if(!defined(__HIP_DEVICE_COMPILE__) || (defined(__gfx9__) && !defined(__gfx906__) && !defined(__gfx900__))

The above should enable compilation for all gfx9 parts except Vega10/20.

BTW it's a shame that compiler doesn't have the __HIP_ARCH_* feature macro that indicates the presence of xDLOPs (MAI) on target, see https://rocm.docs.amd.com/projects/HIP/en/latest/how-to/hip_porting_guide.html#table-of-architecture-properties. Perhaps it's worth requesting it.

@junliume Maybe it is worth introducing intermediate macros that would enable cleaner programming, e.g. put something like this into a header (target_features.hpp?):

#define TARGET_HOST 1 #define TARGET_GFX9 0 #define TARGET_GFX900 0 #define TARGET_GFX906 0 #define TARGET_GFX908 0 #define TARGET_GFX90A 0 #define TARGET_GFX941 0 #define TARGET_GFX942 0 ... #if defined(__HIP_DEVICE_COMPILE__) #undef TARGET_HOST #define TARGET_HOST 0 #if defined(__gfx9__) #undef TARGET_GFX9 #define TARGET_GFX9 1 #endif ... #if defined(__gfx906__) #undef TARGET_GFX906 #define TARGET_GFX906 1 #endif #if defined(__gfx908__) #undef TARGET_GFX908 #define TARGET_GFX908 1 #endif ... #if defined(__gfx940__) #undef TARGET_GFX940 #define TARGET_GFX940 1 #endif #if defined(__gfx941__) #undef TARGET_GFX941 #define TARGET_GFX941 1 #endif ... #endif // Now define features: #define TARGET_GFX94X (TARGET_GFX941 || TARGET_GFX942 || /* enumerate all gfx94s */ ) #define TARGET_GFX_XDLOPS (TARGET_GFX908 || TARGET_GFX90A || TARGET_GFX94X /* add others in the future*/ )

Then you can use this like:

#include <target_features.hpp> #if TARGET_HOST || (TARGET_XDLOPS && TARGET_GFX9) // Enabled for host and for GFX9 XDLOPS devices...

which seems much cleaner to me.

@atamazov found these lines so maybe we should use existing ones there
https://github.com/ROCm/composable_kernel/blob/fix_1371/include/ck/ck.hpp#L55-L74

@junliume At first glance, yes. And you may even find some suitable "feature" macros there, like CK_USE_AMD_MFMA_BF16_1K_OP or so (but I am still unable to find things like TARGET_GFX_XDLOPS there)))

However, these just define, which enforces to use defined() everywhere. What I am proposing is always defining to either 0 or 1, which allows for cleaner constructs. Just to mention.

@junliume At first glance, yes. And you may even find some suitable "feature" macros there, like CK_USE_AMD_MFMA_BF16_1K_OP or so (but I am still unable to find things like TARGET_GFX_XDLOPS there)))

However, these just define, which enforces to use defined() everywhere. What I am proposing is always defining to either 0 or 1, which allows for cleaner constructs. Just to mention.

Good point, I will keep a note about this refactoring and assign CK developer on it.

carlushuang

LGTM

junliume added 2 commits July 2, 2024 18:44

Protect instructions which are only available for gfx94x

ac9a127

remove smfmac tests from unsupported test targets

5d7d367

junliume requested review from zjing14, illsilin, carlushuang and aosewski as code owners July 3, 2024 01:53

junliume added the bug Something isn't working label Jul 3, 2024

fix test name duplicate issue with multiple targets

56b7bae

junliume mentioned this pull request Jul 3, 2024

Limit specific instruction tests to proper platforms #1371

Open

carlushuang previously approved these changes Jul 3, 2024

View reviewed changes

junliume changed the title ~~Remove smfmac tests from unsupported test targets~~ Fix issue with multiple targets and remove smfmac tests from unsupported test targets Jul 3, 2024

junliume mentioned this pull request Jul 3, 2024

fix build issues on gfx1100 alone #1358

Closed

junliume commented Jul 3, 2024

View reviewed changes

Set multi variaty of targets to fix compatibility problem

4b81c7a

junliume dismissed carlushuang’s stale review via 4b81c7a July 3, 2024 03:35

junliume commented Jul 3, 2024

View reviewed changes

junliume added 4 commits July 2, 2024 21:09

fix missing set up arg

0164893

Fix more issue when compiling on gfx1100 and gfx9 at the same time

e6b6e77

fix more issues when tests built are not applicable

6b92926

fix a missing include file

2b5cc1c

aosewski reviewed Jul 3, 2024

View reviewed changes

include/ck/utility/amd_smfmac.hpp Outdated Show resolved Hide resolved

junliume added 2 commits July 3, 2024 00:29

remove duplicated mscro

08d340c

run wmma related tests on proper platforms only

1f86285

jakpiase self-requested a review July 3, 2024 07:57

jakpiase reviewed Jul 3, 2024

View reviewed changes

jakpiase and others added 2 commits July 3, 2024 14:17

disable smfmac tests for ROCm<6.2

81b47b1

GPU target prefer diversity over variations to save time/space

440dd45

junliume added 3 commits July 3, 2024 11:56

test/example applicability should be checked at runtime and not build…

7ba3488

… time

Fix error of consider applicability

b5b429e

These examples should not build when gfx1 in mixed in multi-targets

508b42a

junliume commented Jul 3, 2024

View reviewed changes

client_example/25_wrapper/CMakeLists.txt Outdated Show resolved Hide resolved

Compile time dispatch among multi targets

8ad46ed

junliume added the urgency_blocker blocking feature deliverables label Jul 3, 2024

atamazov reviewed Jul 3, 2024

View reviewed changes

Siplify macros when building multi targets

a7bf0e5

junliume added the CI - Pass label Jul 4, 2024

carlushuang approved these changes Jul 4, 2024

View reviewed changes

junliume merged commit 9590738 into develop Jul 4, 2024
26 checks passed

atamazov mentioned this pull request Jul 23, 2024

[HIP][OpenCL] Do not pass GPU type from the host code, use predefined macros provided by hip-clang instead ROCm/MIOpen#3139

Open

illsilin deleted the fix_1371 branch August 7, 2024 15:30

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Fix issue with multiple targets and remove smfmac tests from unsupported test targets #1372

Fix issue with multiple targets and remove smfmac tests from unsupported test targets #1372

junliume commented Jul 3, 2024

junliume commented Jul 3, 2024

carlushuang left a comment

junliume Jul 3, 2024

junliume Jul 3, 2024

jakpiase left a comment

junliume commented Jul 3, 2024

junliume commented Jul 3, 2024

atamazov Jul 3, 2024 •

edited

Loading

atamazov Jul 3, 2024 •

edited

Loading

junliume Jul 3, 2024

atamazov Jul 4, 2024

junliume Jul 4, 2024

carlushuang left a comment

		#if(!defined(__HIP_DEVICE_COMPILE__) \|\| defined(__gfx908__) \|\| defined(__gfx90a__) \|\| \
		defined(__gfx94__))

Fix issue with multiple targets and remove smfmac tests from unsupported test targets #1372

Fix issue with multiple targets and remove smfmac tests from unsupported test targets #1372

Conversation

junliume commented Jul 3, 2024

junliume commented Jul 3, 2024

carlushuang left a comment

Choose a reason for hiding this comment

junliume Jul 3, 2024

Choose a reason for hiding this comment

junliume Jul 3, 2024

Choose a reason for hiding this comment

jakpiase left a comment

Choose a reason for hiding this comment

junliume commented Jul 3, 2024

junliume commented Jul 3, 2024

atamazov Jul 3, 2024 • edited Loading

Choose a reason for hiding this comment

atamazov Jul 3, 2024 • edited Loading

Choose a reason for hiding this comment

junliume Jul 3, 2024

Choose a reason for hiding this comment

atamazov Jul 4, 2024

Choose a reason for hiding this comment

junliume Jul 4, 2024

Choose a reason for hiding this comment

carlushuang left a comment

Choose a reason for hiding this comment

atamazov Jul 3, 2024 •

edited

Loading

atamazov Jul 3, 2024 •

edited

Loading