Skip to content

Commit 486cdfb

Browse files
authored
Eliminate fp8 gemm on gfx950 (#3903)
FP8 output type not supported by hipblaslt yet on gfx950.
1 parent e244efd commit 486cdfb

File tree

2 files changed

+10
-2
lines changed

2 files changed

+10
-2
lines changed

src/targets/gpu/target.cpp

Lines changed: 8 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -166,6 +166,14 @@ std::vector<pass> target::get_passes(migraphx::context& gctx, const compile_opti
166166
unsupported_fp8ocp_ops.insert("argmax");
167167
unsupported_fp8ocp_ops.insert("argmin");
168168

169+
// disable fp8 dot operations on gfx950
170+
// hipblaslt does not support fp8 gemm with fp8 output yet
171+
const auto device_name = trim(split_string(gpu::get_device_name(), ':').front());
172+
if(starts_with(device_name, "gfx950"))
173+
{
174+
unsupported_fp8ocp_ops.insert("dot");
175+
}
176+
169177
// clang-format off
170178
return
171179
{

test/verify/run_verify.cpp

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -1,7 +1,7 @@
11
/*
22
* The MIT License (MIT)
33
*
4-
* Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved.
4+
* Copyright (c) 2015-2025 Advanced Micro Devices, Inc. All rights reserved.
55
*
66
* Permission is hereby granted, free of charge, to any person obtaining a copy
77
* of this software and associated documentation files (the "Software"), to deal
@@ -32,8 +32,8 @@
3232
#include <migraphx/load_save.hpp>
3333
#include <migraphx/tmp_dir.hpp>
3434
#include <migraphx/verify_args.hpp>
35-
#include <set>
3635

36+
#include <set>
3737
#include <future>
3838
#include <thread>
3939
#include <utility>

0 commit comments

Comments
 (0)