KDB updates for latest assembler #2996

cderb · 2024-05-24T03:08:09Z

This PR updates db_sync and KDB files to support assembler changes and changes made in #2891
Kernel argument changes to LoadProgram function replicated in db_sync test.

junliume · 2024-05-24T19:03:18Z

manually verified fixes of cases in #2891 (comment)

CAHEK7 · 2024-05-25T17:33:37Z

test/gtest/db_sync.cpp

+                    if(miopen::EndsWith(kern.kernel_file, ".s"))
+                    {
+                        compile_options +=
+                            " -mcpu=" +
+                            miopen::LcOptionTargetStrings{handle.GetTargetProperties()}.targetId;
+                    }
+                    else
+                    {
+                        compile_options += " -mcpu=" + handle.GetDeviceName();
+                    }


Suggested change

if(miopen::EndsWith(kern.kernel_file, ".s"))

{

compile_options +=

" -mcpu=" +

miopen::LcOptionTargetStrings{handle.GetTargetProperties()}.targetId;

}

else

{

compile_options += " -mcpu=" + handle.GetDeviceName();

}

compile_options += " -mcpu=";

if(miopen::EndsWith(kern.kernel_file, ".s"))

{

compile_options +=

miopen::LcOptionTargetStrings{handle.GetTargetProperties()}.targetId;

}

else

{

compile_options += handle.GetDeviceName();

}

[R] For me it looks a bit better.

CAHEK7 · 2024-05-25T17:35:02Z

test/gtest/db_sync.cpp

+                            if(miopen::EndsWith(kern.kernel_file, ".s"))
+                            {
+                                compile_options +=
+                                    " -mcpu=" +
+                                    miopen::LcOptionTargetStrings{handle.GetTargetProperties()}
+                                        .targetId;
+                            }
+                            else
+                            {
+                                compile_options += " -mcpu=" + handle.GetDeviceName();
+                            }


Suggested change

if(miopen::EndsWith(kern.kernel_file, ".s"))

{

compile_options +=

" -mcpu=" +

miopen::LcOptionTargetStrings{handle.GetTargetProperties()}

.targetId;

}

else

{

compile_options += " -mcpu=" + handle.GetDeviceName();

}

compile_options += " -mcpu="

if(miopen::EndsWith(kern.kernel_file, ".s"))

{

compile_options +=

miopen::LcOptionTargetStrings{handle.GetTargetProperties()}

.targetId;

}

else

{

compile_options += handle.GetDeviceName();

}

junliume · 2024-05-26T04:52:51Z

@cderb there are quite a few entries not found for gfx90a, could you help to double check? Thanks!

cderb · 2024-05-26T18:25:20Z

@cderb there are quite a few entries not found for gfx90a, could you help to double check? Thanks!

@junliume @JehandadKhan this is a side effect of the targetid being added to the-mcpu argument.
The kernels were generated on a machine with '-mcpu=gfx90a:xnack-'
While the ci machine appears to be using '-mcpu=gfx90a:sramecc+:xnack-'
This of course causes a miss when looking up these arguments.

junliume · 2024-05-27T05:18:26Z

@atamazov can we revert this change?

targetid being added to the-mcpu argument

atamazov · 2024-05-27T17:23:18Z

@junliume

@atamazov can we revert this change?

targetid being added to the-mcpu argument

I do not think so. More info here: #2309 (comment).

I am investigating what needs to be done.

atamazov

🟢 LGTM

atamazov · 2024-05-27T19:38:11Z

~~However, I highly recommend getting back to using different databases for targets with different xnack and sramecc features. Otherwise we potentially losing performance.~~ This is already partially implemented, KDB names include target features, see #2157 (comment). How the kernels are compiled for KDB is another question...

See also #2891 (comment)

Note that we don't need to generate DBs for all possible combinations. We need to do this only for the most common combinations.

cderb · 2024-05-28T15:26:40Z

@atamazov can we revert this change?

targetid being added to the-mcpu argument

@junliume if the kernel arguments change is reverted, then the code changes in this PR are no longer necessary. The kdb changes would also be inaccurate. If the assembler changes are causing an issue, then I can generate the kdbs once more. Otherwise we can throw out this PR.

junliume · 2024-05-28T16:43:35Z

@junliume

@atamazov can we revert this change?

targetid being added to the-mcpu argument

I do not think so. More info here: #2309 (comment).

I am investigating what needs to be done.

@atamazov actually I have the opposite impression, since we build kernel caches using offline compiler (not hipRTC), then the code generated without sramecc or xnack should work in a more general way: e.g. offline compiler generated kernel caches without sramecc or xnack specifics, like "gfx90a" alone, should work on gfx90a regardless of how the GPU reports these detailed features.

update: discussed #2891 (comment)

@atamazov @JehandadKhan please let me know if the above impression is mistaken. Thanks!

If a GPU reports XNACK feature then the code generated without xnack will likely work incorrectly.
If a GPU does not report XNACK then the code generated with xnack may work slower than it potentially can.
If a GPU reports SRAMECC feature then the code generated without sramecc will likely work incorrectly.
If a GPU does not report SRAMECC then the code generated with sramecc may work slower than it potentially can.

cderb · 2024-05-28T16:56:19Z

Confirmed that dbsync will pass with no kdb changes if we revert this change:

MIOpen/src/hip/handlehip.cpp

Lines 500 to 503 in 4f5ed42

    
           else if(miopen::EndsWith(program_name, ".s")) 
        
           { 
        
               params += " -mcpu=" + LcOptionTargetStrings{this->GetTargetProperties()}.targetId; 
        
           }

atamazov · 2024-05-28T22:15:11Z

#2996 (comment) is updated (clarified)

atamazov · 2024-05-29T14:43:00Z

@junliume

since we build kernel caches using offline compiler (not hipRTC),...

The reason is different, I guess. Offline compiler supports xnack target feature at least (well, there is some space for improvement wrt sramecc). But when Tuna builds kernels for KDBs, it uses some special hackery (MIOPEN_DEVICE_ARCH) to substitute actual GPU type and other target info with anything that Tuna wants. This allows to build kernels even on GPUless machines.

then the code generated without sramecc or xnack should work in a more general way: e.g. offline compiler generated kernel caches without sramecc or xnack specifics, like "gfx90a" alone, should work on gfx90a regardless of how the GPU reports these detailed features

Well, actually, if xnack is not specified but the GPU has target feature. then compiler assumes xnack+. And this is not free from the performance POV. I do not know the actual impact for our kernels, but it seems like current system KDB has some space for improvement.

The same considerations apply to sramecc.

If we want to get the best performance for certain targets (I believe we do want this), then we need to

identify the exact list of targets that we need to produce system KDBs
update Tuna and related stuff to produce KDBs for these targets

IIRC the full list of potential targets for KDB is like this:

gfx900
gfx906:sramecc+:xnack-
gfx906:sramecc+:xnack+
-----------------------------------------------------
gfx908:sramecc+:xnack-
gfx908:sramecc+:xnack+
-----------------------------------------------------
gfx90a:sramecc-:xnack-
gfx90a:sramecc-:xnack+
gfx90a:sramecc+:xnack-
gfx90a:sramecc+:xnack+
-----------------------------------------------------
gfx940/941/942:sramecc-:xnack-
gfx940/941/942:sramecc-:xnack+
gfx940/941/942:sramecc+:xnack-
gfx940/941/942:sramecc+:xnack+
-----------------------------------------------------
gfx1030
gfx1031
gfx1100
gfx1101
gfx1102

Again, we don't need to generate KDBs for all possible combinations, only for the important ones.

/cc @junliume @cderb @JehandadKhan

atamazov · 2024-05-29T14:54:35Z

@junliume I am going to open a ticket about "KDB optimization" soon.

cderb · 2024-05-30T17:16:27Z

Discarding this change.

cderb and others added 6 commits May 22, 2024 00:24

kdb update + dbsync update to match PR 2891

121b317

update mi200 kdb, correct db_sync

5060b29

Merge remote-tracking branch 'origin/develop' into cderb/dbsync_argfix

7fbd42a

mi300 kdb

7544e5a

fix lfs

c7eb0e2

mi100 update

99168d0

cderb requested review from JehandadKhan and junliume as code owners May 24, 2024 03:08

cderb and others added 2 commits May 23, 2024 22:08

Merge branch 'develop' into cderb/dbsync_argfix

4880da5

fix clang format issue

6765bfb

junliume added the urgency_blocker label May 24, 2024

junliume and others added 2 commits May 24, 2024 19:14

Merge branch 'develop' into cderb/dbsync_argfix

82247fb

Merge branch 'develop' into cderb/dbsync_argfix

d4852fb

CAHEK7 reviewed May 25, 2024

View reviewed changes

CAHEK7 approved these changes May 25, 2024

View reviewed changes

Merge branch 'develop' into cderb/dbsync_argfix

1e11fbf

atamazov approved these changes May 27, 2024

View reviewed changes

Merge branch 'develop' into cderb/dbsync_argfix

871838d

junliume mentioned this pull request May 28, 2024

Adjustments for the latest assembler (e.g. latest changes in the upstream clang) #2891

Merged

junliume mentioned this pull request May 29, 2024

[WORKAROUND] Disable the fix for #2851 implemented in #2891 #3002

Merged

junliume added the ON_HOLD label May 29, 2024

CAHEK7 mentioned this pull request May 29, 2024

[MHA] turn on 'negative zero as nan' for FP8 numbers #3003

Merged

junliume removed the urgency_blocker label May 29, 2024

cderb closed this May 30, 2024

junliume deleted the cderb/dbsync_argfix branch June 25, 2024 05:07

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

KDB updates for latest assembler #2996

KDB updates for latest assembler #2996

cderb commented May 24, 2024

junliume commented May 24, 2024

CAHEK7 May 25, 2024

CAHEK7 May 25, 2024

junliume commented May 26, 2024

cderb commented May 26, 2024

junliume commented May 27, 2024

atamazov commented May 27, 2024

atamazov left a comment

atamazov commented May 27, 2024 •

edited

Loading

cderb commented May 28, 2024

junliume commented May 28, 2024 •

edited

Loading

cderb commented May 28, 2024

atamazov commented May 28, 2024 •

edited

Loading

atamazov commented May 29, 2024 •

edited

Loading

IIRC the full list of potential targets for KDB is like this:

atamazov commented May 29, 2024

cderb commented May 30, 2024

KDB updates for latest assembler #2996

KDB updates for latest assembler #2996

Conversation

cderb commented May 24, 2024

junliume commented May 24, 2024

CAHEK7 May 25, 2024

Choose a reason for hiding this comment

CAHEK7 May 25, 2024

Choose a reason for hiding this comment

junliume commented May 26, 2024

cderb commented May 26, 2024

junliume commented May 27, 2024

atamazov commented May 27, 2024

atamazov left a comment

Choose a reason for hiding this comment

atamazov commented May 27, 2024 • edited Loading

cderb commented May 28, 2024

junliume commented May 28, 2024 • edited Loading

cderb commented May 28, 2024

atamazov commented May 28, 2024 • edited Loading

atamazov commented May 29, 2024 • edited Loading

IIRC the full list of potential targets for KDB is like this:

atamazov commented May 29, 2024

cderb commented May 30, 2024

atamazov commented May 27, 2024 •

edited

Loading

junliume commented May 28, 2024 •

edited

Loading

atamazov commented May 28, 2024 •

edited

Loading

atamazov commented May 29, 2024 •

edited

Loading