Skip to content
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

Fix sample CI errors #1631

Merged
merged 2 commits into from
Feb 11, 2025
Merged

Fix sample CI errors #1631

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
0f9f57a
Don't treat "no solution" as error in client tools
jichangjichang Feb 7, 2025
9c02cc7
Fix CI errors: Remove F8 Compability checking from hipblaslt runtime
jichangjichang Feb 10, 2025
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
18 changes: 16 additions & 2 deletions clients/benchmarks/client_groupedgemm_fixed_mk.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -2,7 +2,7 @@
*
* MIT License
*
* Copyright (C) 2022-2024 Advanced Micro Devices, Inc.
* Copyright (C) 2022-2025 Advanced Micro Devices, Inc.
*
* Permission is hereby granted, free of charge, to any person obtaining a copy
* of this software and associated documentation files (the "Software"), to deal
Expand Down Expand Up @@ -895,7 +895,21 @@ int test_hipblaslt(hipDataType in_datatype,
if(validIdx.empty())
{
std::cerr << "No Solution found!" << std::endl;
return EXIT_FAILURE;
CHECK_HIP_ERROR(hipFree(d_n_vec));
CHECK_HIP_ERROR(hipFree(d_workspace));
CHECK_HIPBLASLT_ERROR(hipblasLtDestroy(handle));

for(int i = 0; i < gemm_count; i++)
{
CHECK_HIP_ERROR(hipFree(da[i]));
CHECK_HIP_ERROR(hipFree(db[i]));
CHECK_HIP_ERROR(hipFree(dc[i]));
CHECK_HIP_ERROR(hipFree(dd[i]));
if(enable_bias[i])
CHECK_HIP_ERROR(hipFree(d_bias[i]));
}
CHECK_HIP_ERROR(hipStreamDestroy(stream));
return EXIT_SUCCESS;
}

// step2: get default DeviceUserArguments from grouped gemm objects
Expand Down
89 changes: 31 additions & 58 deletions library/src/amd_detail/rocblaslt/src/tensile_host.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -2,7 +2,7 @@
*
* MIT License
*
* Copyright (C) 2022-2024 Advanced Micro Devices, Inc.
* Copyright (C) 2022-2025 Advanced Micro Devices, Inc.
*
* Permission is hereby granted, free of charge, to any person obtaining a copy
* of this software and associated documentation files (the "Software"), to deal
Expand Down Expand Up @@ -314,11 +314,13 @@ namespace
default:;
}

if(typeA == TensileLite::DataType::Float8_fnuz && typeB == TensileLite::DataType::BFloat8_fnuz)
if(typeA == TensileLite::DataType::Float8_fnuz
&& typeB == TensileLite::DataType::BFloat8_fnuz)
{
return TensileLite::DataType::Float8BFloat8_fnuz;
}
else if(typeA == TensileLite::DataType::BFloat8_fnuz && typeB == TensileLite::DataType::Float8_fnuz)
else if(typeA == TensileLite::DataType::BFloat8_fnuz
&& typeB == TensileLite::DataType::Float8_fnuz)
{
return TensileLite::DataType::BFloat8Float8_fnuz;
}
Expand Down Expand Up @@ -422,7 +424,8 @@ namespace
return "f32_bf16_r";
}
else if(typeComputeInput == TensileLite::DataType::Half
&& (typeA == TensileLite::DataType::Float8_fnuz && typeB == TensileLite::DataType::Half
&& (typeA == TensileLite::DataType::Float8_fnuz
&& typeB == TensileLite::DataType::Half
|| typeA == TensileLite::DataType::Half
&& typeB == TensileLite::DataType::Float8_fnuz))
{
Expand Down Expand Up @@ -464,7 +467,8 @@ namespace
return "c_f32_fast_bf16_r";
}
else if(typeComputeInput == TensileLite::DataType::Half
&& (typeA == TensileLite::DataType::Float8_fnuz && typeB == TensileLite::DataType::Half
&& (typeA == TensileLite::DataType::Float8_fnuz
&& typeB == TensileLite::DataType::Half
|| typeA == TensileLite::DataType::Half
&& typeB == TensileLite::DataType::Float8_fnuz))
{
Expand Down Expand Up @@ -1066,7 +1070,8 @@ namespace
auto k = prob.k && alpha ? prob.k : 0;

// fallback to f32 for f16 compute type after alpha/beta assignment
if (prob.compute_type == rocblaslt_compute_f16) {
if(prob.compute_type == rocblaslt_compute_f16)
{
compute_type = roc2TensileType(prob.compute_type);
}

Expand Down Expand Up @@ -1331,7 +1336,8 @@ namespace
assignAlphaBeta(compute_type, prob.alpha, prob.beta, &alpha, &beta);

// fallback to f32 for f16 compute type after alpha/beta assignment
if (prob.compute_type == rocblaslt_compute_f16) {
if(prob.compute_type == rocblaslt_compute_f16)
{
compute_type = roc2TensileType(prob.compute_type);
}

Expand Down Expand Up @@ -1471,27 +1477,31 @@ namespace
{TensileLite::DataType::Double, (double)0.0},
};

if (argument_vals.find(compute_type) == argument_vals.end()) {
if(argument_vals.find(compute_type) == argument_vals.end())
{
log_error(__func__, "Unsupported compute type");
throw std::runtime_error("[GetTensileInputs] unsupported compute type.");
}

// push 2 activation arguments
std::visit([&inputs, &prob](auto val) {
inputs.activationArgs.push_back(val);
inputs.activationArgs.push_back(val);
if(prob.k)
inputs.alpha = *(decltype(val)*)(prob.alpha);
else
inputs.alpha = val;
inputs.beta = *(decltype(val)*)(prob.beta);
}, argument_vals.at(compute_type));
std::visit(
[&inputs, &prob](auto val) {
inputs.activationArgs.push_back(val);
inputs.activationArgs.push_back(val);
if(prob.k)
inputs.alpha = *(decltype(val)*)(prob.alpha);
else
inputs.alpha = val;
inputs.beta = *(decltype(val)*)(prob.beta);
},
argument_vals.at(compute_type));

// convert alpha and beta to float if compute type is half
if (prob.compute_type == rocblaslt_compute_f16) {
if(prob.compute_type == rocblaslt_compute_f16)
{
inputs.activationArgs = {0.0f, 0.0f};
inputs.alpha = static_cast<float>(std::get<hipblasLtHalf>(inputs.alpha));
inputs.beta = static_cast<float>(std::get<hipblasLtHalf>(inputs.beta));
inputs.alpha = static_cast<float>(std::get<hipblasLtHalf>(inputs.alpha));
inputs.beta = static_cast<float>(std::get<hipblasLtHalf>(inputs.beta));
}

return inputs;
Expand Down Expand Up @@ -1613,7 +1623,7 @@ namespace
}

// TensileHost is not copyable or assignable
TensileHost(const TensileHost&) = delete;
TensileHost(const TensileHost&) = delete;
TensileHost& operator=(const TensileHost&) = delete;

// Get the number of devices
Expand Down Expand Up @@ -2983,38 +2993,6 @@ rocblaslt_status getBestSolutions(RocblasltContractionProblem const& prob,
return rocblaslt_status_success;
}

void checkF8Compatiblity(const std::string &deviceString, const TensileLite::ContractionProblemGemm& prob) {

bool isGFX94X = deviceString.find("gfx940") != std::string::npos ||
deviceString.find("gfx941") != std::string::npos ||
deviceString.find("gfx942") != std::string::npos;

auto isFNUZ = [](TensileLite::DataType type) {
return type == TensileLite::DataType::Float8_fnuz ||
type == TensileLite::DataType::BFloat8_fnuz;
};

auto isOCP = [](TensileLite::DataType type) {
return type == TensileLite::DataType::Float8 ||
type == TensileLite::DataType::BFloat8;
};

bool hasFNUZ = isFNUZ(prob.a().dataType()) ||
isFNUZ(prob.b().dataType()) ||
isFNUZ(prob.c().dataType()) ||
isFNUZ(prob.d().dataType());

bool hasOCP = isOCP(prob.a().dataType()) ||
isOCP(prob.b().dataType()) ||
isOCP(prob.c().dataType()) ||
isOCP(prob.d().dataType());

if((hasFNUZ && !isGFX94X) || (hasOCP && isGFX94X) || (hasFNUZ && hasOCP)) {
log_error(__func__, "Requested F8 type not supported");
throw std::runtime_error("[checkF8] Requested F8 type not supported.");
}
}

template <typename MyProblem>
rocblaslt_status getAllSolutions(MyProblem& prob,
rocblaslt_handle handle,
Expand Down Expand Up @@ -3044,16 +3022,11 @@ rocblaslt_status getAllSolutions(MyProblem&

if constexpr(std::is_same<MyProblem, TensileLite::ContractionProblemGemm>::value)
{
checkF8Compatiblity(deviceString, prob);

solutions = library->findAllSolutions(
prob, *hardware, TensileLite::SolutionLibrarySearchType::GEMM_TYPE_ONLY);
}
else if constexpr(std::is_same<MyProblem, TensileLite::ContractionProblemGroupedGemm>::value)
{
for (const auto &gemm : prob.gemms)
checkF8Compatiblity(deviceString, gemm);

solutions = library->findAllSolutionsGroupedGemm(
prob.gemms, *hardware, TensileLite::SolutionLibrarySearchType::GEMM_TYPE_ONLY);
}
Expand Down