From 1880898812e4b90c1bf7859acf6c57aab03a91d3 Mon Sep 17 00:00:00 2001 From: tvukovic-amd Date: Fri, 17 May 2024 08:42:12 +0100 Subject: [PATCH 1/4] Fixed rocblas failed tests --- src/targets/gpu/target.cpp | 4 ++++ test/gpu/gemm_tune.cpp | 2 +- 2 files changed, 5 insertions(+), 1 deletion(-) diff --git a/src/targets/gpu/target.cpp b/src/targets/gpu/target.cpp index 3264f298f62..5e71c29a3d4 100644 --- a/src/targets/gpu/target.cpp +++ b/src/targets/gpu/target.cpp @@ -103,6 +103,10 @@ std::vector target::get_passes(migraphx::context& gctx, const compile_opti unsupported_fp8_ops.insert("dot"); unsupported_fp8_ops.insert("quant_dot"); } +#else + // mlir doesn't support fp8 dot + unsupported_fp8_ops.insert("dot"); + unsupported_fp8_ops.insert("quant_dot"); #endif // MIOpen doesn't have support for fp8 pooling yet. unsupported_fp8_ops.insert("pooling"); diff --git a/test/gpu/gemm_tune.cpp b/test/gpu/gemm_tune.cpp index 75081bc37ba..52a32a771b2 100644 --- a/test/gpu/gemm_tune.cpp +++ b/test/gpu/gemm_tune.cpp @@ -182,7 +182,6 @@ TEST_CASE(gemm_tune_strided_lowered) EXPECT(0 == solution_idx.to()); #endif } -#endif TEST_CASE(gemm_tune_invalid_sol_index) { @@ -223,5 +222,6 @@ TEST_CASE(gemm_tune_invalid_sol_index) EXPECT(0 != solution_idx.to()); #endif } +#endif int main(int argc, const char* argv[]) { test::run(argc, argv); } From 1b19621b08dd89b81184539805bf3be418d4ddfd Mon Sep 17 00:00:00 2001 From: tvukovic-amd Date: Fri, 17 May 2024 10:39:30 +0100 Subject: [PATCH 2/4] Fix format --- src/targets/gpu/target.cpp | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/src/targets/gpu/target.cpp b/src/targets/gpu/target.cpp index 5e71c29a3d4..58dbdd7f50c 100644 --- a/src/targets/gpu/target.cpp +++ b/src/targets/gpu/target.cpp @@ -104,9 +104,9 @@ std::vector target::get_passes(migraphx::context& gctx, const compile_opti unsupported_fp8_ops.insert("quant_dot"); } #else - // mlir doesn't support fp8 dot - unsupported_fp8_ops.insert("dot"); - unsupported_fp8_ops.insert("quant_dot"); + // mlir doesn't support fp8 dot + unsupported_fp8_ops.insert("dot"); + unsupported_fp8_ops.insert("quant_dot"); #endif // MIOpen doesn't have support for fp8 pooling yet. unsupported_fp8_ops.insert("pooling"); From eb15631c2f2e6348d06eb5be4f68955e37d55ee1 Mon Sep 17 00:00:00 2001 From: tvukovic-amd Date: Mon, 20 May 2024 09:54:37 +0100 Subject: [PATCH 3/4] Add changes according to comments --- .../gpu/include/migraphx/gpu/rocblas.hpp | 4 ++-- src/targets/gpu/rocblas.cpp | 18 +++++++++++------- src/targets/gpu/target.cpp | 6 ------ 3 files changed, 13 insertions(+), 15 deletions(-) diff --git a/src/targets/gpu/include/migraphx/gpu/rocblas.hpp b/src/targets/gpu/include/migraphx/gpu/rocblas.hpp index a03fa8385d2..d23c40f9d05 100644 --- a/src/targets/gpu/include/migraphx/gpu/rocblas.hpp +++ b/src/targets/gpu/include/migraphx/gpu/rocblas.hpp @@ -38,13 +38,13 @@ using rocblas_handle_ptr = MIGRAPHX_MANAGE_PTR(rocblas_handle, rocblas_destroy_h rocblas_handle_ptr create_rocblas_handle_ptr(); rocblas_handle_ptr create_rocblas_handle_ptr(hipStream_t s); - +#endif struct context; MIGRAPHX_GPU_EXPORT bool get_compute_fp32_flag(); MIGRAPHX_GPU_EXPORT bool rocblas_fp8_available(); -#endif + } // namespace gpu } // namespace MIGRAPHX_INLINE_NS } // namespace migraphx diff --git a/src/targets/gpu/rocblas.cpp b/src/targets/gpu/rocblas.cpp index 5779f7ae6e1..6bb045d8fd9 100644 --- a/src/targets/gpu/rocblas.cpp +++ b/src/targets/gpu/rocblas.cpp @@ -48,7 +48,7 @@ rocblas_handle_ptr create_rocblas_handle_ptr(hipStream_t s) rocblas_set_stream(rb.get(), s); return rb; } - +#endif bool get_compute_fp32_flag() { const auto device_name = trim(split_string(get_device_name(), ':').front()); @@ -57,13 +57,17 @@ bool get_compute_fp32_flag() bool rocblas_fp8_available() { -#ifndef MIGRAPHX_USE_ROCBLAS_FP8_API - return false; -#else - return gfx_has_fp8_intrinsics(); -#endif + #if MIGRAPHX_USE_ROCBLAS + #ifndef MIGRAPHX_USE_ROCBLAS_FP8_API + return false; + #else + return gfx_has_fp8_intrinsics(); + #endif + #else + return false; + #endif } -#endif + } // namespace gpu } // namespace MIGRAPHX_INLINE_NS } // namespace migraphx diff --git a/src/targets/gpu/target.cpp b/src/targets/gpu/target.cpp index 58dbdd7f50c..4a18e25aab5 100644 --- a/src/targets/gpu/target.cpp +++ b/src/targets/gpu/target.cpp @@ -97,17 +97,11 @@ std::vector target::get_passes(migraphx::context& gctx, const compile_opti unsupported_types.erase(shape::type_t::tuple_type); // whiltelist supported Ops for the FP8 std::set unsupported_fp8_ops = {}; -#if MIGRAPHX_USE_ROCBLAS if(not gpu::rocblas_fp8_available()) { unsupported_fp8_ops.insert("dot"); unsupported_fp8_ops.insert("quant_dot"); } -#else - // mlir doesn't support fp8 dot - unsupported_fp8_ops.insert("dot"); - unsupported_fp8_ops.insert("quant_dot"); -#endif // MIOpen doesn't have support for fp8 pooling yet. unsupported_fp8_ops.insert("pooling"); if(not gpu::gfx_has_fp8_intrinsics()) From dfb745031661d409b585cb962f2e4fe1008fb6f8 Mon Sep 17 00:00:00 2001 From: tvukovic-amd Date: Mon, 20 May 2024 12:26:29 +0100 Subject: [PATCH 4/4] Fix licencing and format --- src/targets/gpu/rocblas.cpp | 18 +++++++++--------- test/CMakeLists.txt | 2 +- 2 files changed, 10 insertions(+), 10 deletions(-) diff --git a/src/targets/gpu/rocblas.cpp b/src/targets/gpu/rocblas.cpp index 6bb045d8fd9..798fefbb811 100644 --- a/src/targets/gpu/rocblas.cpp +++ b/src/targets/gpu/rocblas.cpp @@ -57,15 +57,15 @@ bool get_compute_fp32_flag() bool rocblas_fp8_available() { - #if MIGRAPHX_USE_ROCBLAS - #ifndef MIGRAPHX_USE_ROCBLAS_FP8_API - return false; - #else - return gfx_has_fp8_intrinsics(); - #endif - #else - return false; - #endif +#if MIGRAPHX_USE_ROCBLAS +#ifndef MIGRAPHX_USE_ROCBLAS_FP8_API + return false; +#else + return gfx_has_fp8_intrinsics(); +#endif +#else + return false; +#endif } } // namespace gpu diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt index b6ecc2e35d3..80edc6121ab 100644 --- a/test/CMakeLists.txt +++ b/test/CMakeLists.txt @@ -1,7 +1,7 @@ # #################################################################################### # The MIT License (MIT) # -# Copyright (c) 2015-2023 Advanced Micro Devices, Inc. All rights reserved. +# Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. # # Permission is hereby granted, free of charge, to any person obtaining a copy # of this software and associated documentation files (the "Software"), to deal