Skip to content

[hipBLASLt dev] Refined the CU-Fallback case for XCC value #483

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

Merged
solaslin merged 1 commit into from
Jul 11, 2025
Merged

[hipBLASLt dev] Refined the CU-Fallback case for XCC value #483

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
23 changes: 14 additions & 9 deletions projects/hipblaslt/tensilelite/Tensile/Source/lib/source/AMDGPU.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -28,11 +28,16 @@

namespace TensileLite
{
static const std::map<AMDGPU::Processor, std::vector<int>>& NonStandardCUMap()
// TODD- Currently, this map is for detecting CU-Fallback + runtime modification for XCC.
// XCC is introduced since 942, so this is only targeting at gfx942 for now.
// (i.e. Detect this on 942 only. For other devices, we don't have problems for CU-FB + mismatching XCC,
// so all the other devices will be seen as standardCU)
// THIS IS A TEMPORARILY WORKAROUND. WE MAY NEED TO COME UP WITH A BETTER WAY TO DETECT CU-FALLBACK
static const std::map<AMDGPU::Processor, int>& StandardCUMap()
{
static const std::map<AMDGPU::Processor, std::vector<int>> NonStandardCU = {
{AMDGPU::Processor::gfx90a, {104}}, {AMDGPU::Processor::gfx942, {20, 38, 64, 80, 228}}};
return NonStandardCU;
static const std::map<AMDGPU::Processor, int> StandardCU_XCC
= {{AMDGPU::Processor::gfx942, 304}};
return StandardCU_XCC;
}

TENSILE_API std::string AMDGPU::type() const
Expand Down Expand Up @@ -66,12 +71,12 @@ namespace TensileLite
// assume current device is a standard cu device.
isStandardCUs = 1;

auto mapIter = NonStandardCUMap().find(processor);
if(mapIter != NonStandardCUMap().end())
auto mapIter = StandardCUMap().find(processor);
// NB: For any other devices not included in the map, we always see them as true (they have no XCC predicate issue.)
if(mapIter != StandardCUMap().end())
{
auto list = mapIter->second;
// current device is a speicalized cu device.
if(std::count(list.begin(), list.end(), computeUnitCount) > 0)
// check if current device is a standard cu devcie.
if(computeUnitCount != mapIter->second)
isStandardCUs = 0;
}

Expand Down