[miopen] Log applicability failures by SreecharanGundaboluAMD · Pull Request #3964 · ROCm/rocm-libraries

SreecharanGundaboluAMD · 2026-01-20T02:39:21Z

This pull request introduces a set of macros and standardized messages to improve the clarity, consistency, and maintainability of applicability checks in solver classes across the MIOpen codebase. The main theme is refactoring the IsApplicable methods of various solvers to use the new macros for logging and returning early when a solver is not applicable, along with standardized messages for common reasons. This change helps developers quickly understand why a solver was rejected and ensures uniform logging across different solver implementations.

The most important changes are:

Core Infrastructure:

Added new header file solver_utils.hpp that defines the MIOPEN_SOLVER_INAPPLICABLE_IF and MIOPEN_SOLVER_INAPPLICABLE_IF_CONTEXT macros, as well as a set of standardized inapplicability message strings in the inapplicable_msg namespace. This file includes documentation and usage guidelines for the macros.

Refactoring Applicability Checks (Activ Solvers):

Refactored IsApplicable methods in activ/bwd_0.cpp, activ/bwd_1.cpp, activ/fwd_0.cpp, and activ/fwd_1.cpp to use the new macros for early returns and standardized messages, replacing manual if checks and custom strings. This improves logging and code consistency for activation solvers. [1] [2] [3] [4] [5] [6] [7] [8]

Refactoring Applicability Checks (Adam Solvers):

Updated IsApplicable methods in adam/adam.cpp and adam/transformers_adam_w.cpp to use the new macro and standardized message for contiguous tensor checks. [1] [2]

Refactoring Applicability Checks (BatchNorm Solver):

Refactored IsApplicable in batchnorm/backward_per_activation.cpp to use the macros and standardized messages for 2D checks, direction/mode checks, and data type validation, replacing multiple manual if statements.

Codebase Integration:

Added #include <miopen/solver/solver_utils.hpp> to all affected solver source files to enable use of the new macros and messages. [1] [2] [3] [4] [5] [6] [7]

These changes collectively standardize the way solver applicability is checked and logged, making the code easier to maintain and debug.

Submission Checklist

Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.

…ndaboluAMD/log_applicability_failures

…ility_failures

AnaghaRaoAMD · 2026-01-30T02:38:34Z

-        return false;
+    MIOPEN_SOLVER_INAPPLICABLE_IF(
+        !(x_elem_sz == y_elem_sz && dx_elem_sz == dy_elem_sz && x_elem_sz == dx_elem_sz),
+        "Element sizes do not match");


Can this be added to inapplicable_msg?

or we could use ElementSizeMismatch?

AnaghaRaoAMD · 2026-01-30T03:00:49Z

+inline constexpr const char* Layout            = "Unsupported tensor layout.";
+inline constexpr const char* UnsupportedDevice = "Unsupported GPU.";
+inline constexpr const char* Direction         = "Unsupported direction.";
+inline constexpr const char* UseAsmKernels     = "Assembly kernels are disabled.";


nit: Name seems counter intuitive

Suggested change

inline constexpr const char* UseAsmKernels = "Assembly kernels are disabled.";

inline constexpr const char* AsmKernelsDisabled = "Assembly kernels are disabled.";

AnaghaRaoAMD · 2026-01-30T03:39:49Z

+
+    MIOPEN_SOLVER_INAPPLICABLE_IF(!ctx.rmv.IsV2orV3(), inapplicable_msg::MetaData);
+
+    MIOPEN_SOLVER_INAPPLICABLE_IF(!(problem.IsFp32() || problem.IsFp16()),


missed IsBfp8 type

AnaghaRaoAMD · 2026-01-30T04:06:55Z

-
-    if(problem.IsTensorsCasted() || problem.IsFp8() || problem.IsBfp8())
-        return false;


Missed Datatype check for fp8 and bfp8

AnaghaRaoAMD · 2026-01-30T04:08:20Z

-
-    if(problem.IsTensorsCasted())
-        return false;
+    const std::string solver = SolverDbId();


AnaghaRaoAMD · 2026-01-30T04:12:52Z

+
+    MIOPEN_SOLVER_INAPPLICABLE_IF(!problem.IsLayoutDefault(), inapplicable_msg::Layout);
+
+    MIOPEN_SOLVER_INAPPLICABLE_IF(problem.IsTensorsCasted() || problem.IsFp8() || problem.IsBfp8(),


should be separated to IsTensorsCasted and DataType msg for uniformity

AnaghaRaoAMD · 2026-01-30T04:29:21Z

-        if(current_block_size * current_grid_size * current_splits_4G > 0xffffffffULL)
-            return false;
+        MIOPEN_SOLVER_INAPPLICABLE_IF(
+            (current_block_size * current_grid_size * current_splits_4G > 0xffffffffULL),


can add a new worksize limit exceeded msg

AnaghaRaoAMD · 2026-01-30T04:33:07Z

+        (!problem.IsFp32() && !problem.IsFp16() &&
+         !(problem.IsBfp16() && (device_name == "gfx90a" || device_name == "gfx942" ||
+                                 StartsWith(device_name, "gfx95")))),
+        "Datatype and HW combination not supported");


Suggested change

"Datatype and HW combination not supported");

"Datatype and device combination not supported");

AnaghaRaoAMD · 2026-01-30T04:33:47Z

-        if(current_block_size * current_grid_size * current_splits_4G > 0xffffffffULL)
-            return false;
+        MIOPEN_SOLVER_INAPPLICABLE_IF(
+            (current_block_size * current_grid_size * current_splits_4G > 0xffffffffULL),


can be replace with worksize msg

AnaghaRaoAMD · 2026-01-30T04:35:18Z

+        (!problem.IsFp32() && !problem.IsFp16() &&
+         !(problem.IsBfp16() && (device_name == "gfx90a" || device_name == "gfx942" ||
+                                 StartsWith(device_name, "gfx95")))),
+        "Datatype, HW combination not supported");


Suggested change

"Datatype, HW combination not supported");

"Datatype and device combination not supported");

AnaghaRaoAMD · 2026-01-30T04:35:49Z

-        if(current_block_size * current_grid_size * current_splits_4G > 0xffffffffULL)
-            return false;
+        MIOPEN_SOLVER_INAPPLICABLE_IF(
+            !(current_block_size * current_grid_size * current_splits_4G > 0xffffffffULL),


can be replaced with worksize msg

AnaghaRaoAMD · 2026-01-30T05:02:25Z

+    MIOPEN_SOLVER_INAPPLICABLE_IF(!ctx.use_hip_kernels, inapplicable_msg::HIPDisabled);
+
+    MIOPEN_SOLVER_INAPPLICABLE_IF(!legacy_ck::is_ck_supported_hardware(ctx.GetStream()),
+                                  "Legacy CK does not support HW");


Suggested change

"Legacy CK does not support HW");

"Legacy CK is not support on this Hardware");

HW is misleading for Height and Width parameters

AnaghaRaoAMD · 2026-01-30T05:04:32Z

-        legacy_ck::get_ck_convolution_problem_descriptor(problem));
+    MIOPEN_SOLVER_INAPPLICABLE_IF(!ctx.use_hip_kernels, inapplicable_msg::HIPDisabled);
+
+    MIOPEN_SOLVER_INAPPLICABLE_IF(!legacy_ck::is_ck_supported_hardware(ctx.GetStream()),


Can also use CKWhitelist, "GPU is not in CK whitelist."

AnaghaRaoAMD · 2026-01-30T05:09:20Z

+    MIOPEN_SOLVER_INAPPLICABLE_IF(!problem.AllTensorsLengthsFitIntoInt(),
+                                  inapplicable_msg::AllTensorsDimsFitIntoInt);
+
+    MIOPEN_SOLVER_INAPPLICABLE_IF(!problem.IsLayoutDefault() && !problem.IsLayoutNHWC(),


MIOPEN_SOLVER_INAPPLICABLE_IF(!problem.IsLayoutDefault() && !problem.IsLayoutNHWC(),
added twice

AnaghaRaoAMD · 2026-01-30T05:10:34Z

-            return false;
-        if(test_cast(problem.GetWeights()))
-            return false;
+        MIOPEN_SOLVER_INAPPLICABLE_IF(test_cast(problem.GetIn()), "Input tensor missing cast type");


Suggested change

MIOPEN_SOLVER_INAPPLICABLE_IF(test_cast(problem.GetIn()), "Input tensor missing cast type");

MIOPEN_SOLVER_INAPPLICABLE_IF(test_cast(problem.GetOut()), "Output tensor missing cast type");

AnaghaRaoAMD · 2026-01-30T06:10:49Z


-    if(!ConvDirectNaiveConvIsApplicableByKernelType(ctx, problem))
-        return false;
+    MIOPEN_SOLVER_INAPPLICABLE_IF(!problem.IsLayoutDefault() && !problem.IsLayoutNHWC(),


Suggested change

MIOPEN_SOLVER_INAPPLICABLE_IF(!problem.IsLayoutDefault() && !problem.IsLayoutNHWC(),

MIOPEN_SOLVER_INAPPLICABLE_IF(!ConvDirectNaiveConvIsApplicableByKernelType(ctx, problem),

inapplicable_msg::Generic);

AnaghaRaoAMD · 2026-01-30T06:11:30Z

+    MIOPEN_SOLVER_INAPPLICABLE_IF(!problem.IsLayoutDefault() && !problem.IsLayoutNHWC(),
+                                  inapplicable_msg::Layout);


Suggested change

MIOPEN_SOLVER_INAPPLICABLE_IF(!problem.IsLayoutDefault() && !problem.IsLayoutNHWC(),

inapplicable_msg::Layout);

MIOPEN_SOLVER_INAPPLICABLE_IF(!ConvDirectNaiveConvIsApplicableByKernelType(ctx, problem),

inapplicable_msg::Generic);

AnaghaRaoAMD · 2026-01-30T06:12:34Z

-            return false;
+        MIOPEN_SOLVER_INAPPLICABLE_IF(test_cast(problem.GetIn()), "Input tensor missing cast type");
+        MIOPEN_SOLVER_INAPPLICABLE_IF(test_cast(problem.GetOut()),
+                                      "Weight tensor missing cast type");


Suggested change

"Weight tensor missing cast type");

"Output tensor missing cast type");

AnaghaRaoAMD · 2026-01-30T06:34:56Z

    std::tie(gemm_g, gemm_m, gemm_n, gemm_k_total) = CalculateGemmSize(problem);

-    return static_ck::IsValidGridGemmXdlops(gemm_m, gemm_n, gemm_k_total);
+    MIOPEN_SOLVER_INAPPLICABLE_IF(!static_ck::IsValidGridGemmXdlops(gemm_m, gemm_n, gemm_k_total),


is this a case for InvalidGridGemm

AnaghaRaoAMD · 2026-01-30T06:36:58Z


-    return static_ck::IsValidGridGemmXdlops(gemm_m, gemm_n, gemm_k_total);
+    MIOPEN_SOLVER_INAPPLICABLE_IF(!static_ck::IsValidGridGemmXdlops(gemm_m, gemm_n, gemm_k_total),
+                                  inapplicable_msg::NoKernelForConfig);


Suggested change

inapplicable_msg::NoKernelForConfig);

inapplicable_msg::InvalidGridGemm);

AnaghaRaoAMD · 2026-01-30T06:52:37Z


-    return static_ck::IsValidGridGemmXdlops(gemm_m, gemm_n, gemm_k_total);
+    MIOPEN_SOLVER_INAPPLICABLE_IF(!static_ck::IsValidGridGemmXdlops(gemm_m, gemm_n, gemm_k_total),
+                                  inapplicable_msg::NoKernelForConfig);


Suggested change

inapplicable_msg::NoKernelForConfig);

inapplicable_msg::InvalidGridGemm);

AnaghaRaoAMD · 2026-01-30T07:00:33Z

-        if(IS_DISABLED(MIOPEN_DEBUG_AMD_MP_BD_WINOGRAD_F3X3))
-            return false;
+        MIOPEN_SOLVER_INAPPLICABLE_IF(IS_DISABLED(MIOPEN_DEBUG_AMD_MP_BD_WINOGRAD_F3X3),
+                                      inapplicable_msg::EnvDisabled);


It would be useful to also capture the ENV VAR in EnvDisabled msg to help isolate why its disabled

AnaghaRaoAMD · 2026-01-30T07:05:12Z

+          // #if 0
+          //            // There is a stronger restriction than this one, which make this one
+          //            unnecessary.
+          //            // The kernel read stripes (in height direction, one stripe at a time) of
+          //            input into LDS,
+          //            // the height of stripe is (MLO_N_ALIGNED_OUT_SCAN_BLK - 1) *
+          //            MLO_FILTER_STRIDE1 +
+          //            // MLO_FILTER_SIZE1, (MLO_FILTER_SIZE1 - MLO_FILTER_STRIDE1) of it is
+          //            reusable from
+          //            // previous read, (MLO_N_ALIGNED_OUT_SCAN_BLK * MLO_FILTER_STRIDE1) of it is
+          //            fresh read
+          //            // from device memory. So (MLO_FILTER_SIZE1 - MLO_FILTER_STRIDE1) need no
+          //            less than 0.
+          //            // TODO: chao: revisit this if failure is encountered.
+          //            problem.GetWeightsHeight() >= problem.GetKernelStrideH() &&
+          // #endif


Suggested change

// #if 0

// // There is a stronger restriction than this one, which make this one

// unnecessary.

// // The kernel read stripes (in height direction, one stripe at a time) of

// input into LDS,

// // the height of stripe is (MLO_N_ALIGNED_OUT_SCAN_BLK - 1) *

// MLO_FILTER_STRIDE1 +

// // MLO_FILTER_SIZE1, (MLO_FILTER_SIZE1 - MLO_FILTER_STRIDE1) of it is

// reusable from

// // previous read, (MLO_N_ALIGNED_OUT_SCAN_BLK * MLO_FILTER_STRIDE1) of it is

// fresh read

// // from device memory. So (MLO_FILTER_SIZE1 - MLO_FILTER_STRIDE1) need no

// less than 0.

// // TODO: chao: revisit this if failure is encountered.

// problem.GetWeightsHeight() >= problem.GetKernelStrideH() &&

// #endif

#if 0

// There is a stronger restriction than this one, which make this one unnecessary.

// The kernel read stripes (in height direction, one stripe at a time) of input into LDS,

// the height of stripe is (MLO_N_ALIGNED_OUT_SCAN_BLK - 1) * MLO_FILTER_STRIDE1 +

// MLO_FILTER_SIZE1, (MLO_FILTER_SIZE1 - MLO_FILTER_STRIDE1) of it is reusable from

// previous read, (MLO_N_ALIGNED_OUT_SCAN_BLK * MLO_FILTER_STRIDE1) of it is fresh read

// from device memory. So (MLO_FILTER_SIZE1 - MLO_FILTER_STRIDE1) need no less than 0.

// TODO: chao: revisit this if failure is encountered.

problem.GetWeightsHeight() >= problem.GetKernelStrideH() &&

#endif

AnaghaRaoAMD · 2026-01-30T07:08:42Z

-        if (!supported)
-            return false;
+        MIOPEN_SOLVER_INAPPLICABLE_IF(!supported,
+                                              inapplicable_msg::Workaround);


I am not sure if this is due to an explicit workaround?

AnaghaRaoAMD · 2026-01-30T07:12:12Z

 bool ConvWinoFuryRxSCommon<Winodata, Winofilter>::IsApplicable(const ExecutionContext& ctx,
                                                               const ProblemDescription& problem)
 {
+    std::string context = "ConvWinoFuryRxSCommon"; // for logging


do we mean to add it to all solvers? or this is leftover code?

AnaghaRaoAMD · 2026-01-30T07:14:28Z

-        return false;
-    if(desc.op_map[0]->kind() != miopenFusionOpConvForward)
-        return false;
+    MIOPEN_SOLVER_INAPPLICABLE_IF((desc.op_map.size() > 3), "Too many ops in fusion plan.");


FusionOpCount

AnaghaRaoAMD · 2026-01-30T07:30:05Z

-           !(gemm::IsAnyBufferBf16(xDesc, yDesc, wDesc) && !gemm::IsBf16Supported) &&
-           !(gemm::IsAnyBufferFp16(xDesc, yDesc, wDesc) && !gemm::IsFp16Supported);
+    MIOPEN_SOLVER_INAPPLICABLE_IF((problem.IsFp8() && !rblas_fp8_supported),
+                                  "GEMM not applicable for F8 on this GPU architecture");


Suggested change

"GEMM not applicable for F8 on this GPU architecture");

"GEMM not applicable for FP8 on this GPU architecture");

AnaghaRaoAMD · 2026-01-30T07:31:01Z

-        }
+        MIOPEN_SOLVER_INAPPLICABLE_IF(
+            !rblas_fp8_supported,
+            "GEMM not supported with casted tensors on this GPU architecture");


IsTensorsCasted?

AnaghaRaoAMD · 2026-01-30T07:31:26Z

-           !(gemm::IsAnyBufferBf16(dxDesc, dyDesc, wDesc) && !gemm::IsBf16Supported) &&
-           !(gemm::IsAnyBufferFp16(dxDesc, dyDesc, wDesc) && !gemm::IsFp16Supported);
+    MIOPEN_SOLVER_INAPPLICABLE_IF((problem.IsFp8() && !rblas_fp8_supported),
+                                  "GEMM not applicable for F8 on this GPU architecture");


Suggested change

"GEMM not applicable for F8 on this GPU architecture");

"GEMM not applicable for FP8 on this GPU architecture");

AnaghaRaoAMD · 2026-01-30T07:32:21Z

-        }
+        MIOPEN_SOLVER_INAPPLICABLE_IF(
+            !rblas_fp8_supported,
+            "GEMM not supported with casted tensors on this GPU architecture");


IsTensorsCasted

AnaghaRaoAMD · 2026-01-30T07:33:21Z

-           !(gemm::IsAnyBufferBf16(xDesc, dyDesc, dwDesc) && !gemm::IsBf16Supported) &&
-           !(gemm::IsAnyBufferFp16(xDesc, dyDesc, dwDesc) && !gemm::IsFp16Supported);
+    MIOPEN_SOLVER_INAPPLICABLE_IF((problem.IsFp8() && !rblas_fp8_supported),
+                                  "GEMM not applicable for F8 on this GPU architecture");


Suggested change

"GEMM not applicable for F8 on this GPU architecture");

"GEMM not applicable for FP8 on this GPU architecture");

AnaghaRaoAMD · 2026-01-30T07:35:46Z

-        return false;
+    MIOPEN_SOLVER_INAPPLICABLE_IF(!problem.IsSameType(), inapplicable_msg::DataTypeMismatch);
+
+    MIOPEN_SOLVER_INAPPLICABLE_IF(!IsLargeIndex(problem), "Not a large index problem");


Suggested change

MIOPEN_SOLVER_INAPPLICABLE_IF(!IsLargeIndex(problem), "Not a large index problem");

MIOPEN_SOLVER_INAPPLICABLE_IF(!IsLargeIndex(problem), inapplicable_msg::IndexRange);

AnaghaRaoAMD · 2026-01-30T07:36:57Z

-        return false;
+    MIOPEN_SOLVER_INAPPLICABLE_IF(!problem.IsAllContiguous(), inapplicable_msg::NotContiguous);
+
+    MIOPEN_SOLVER_INAPPLICABLE_IF((problem.GetDim() != 0), "Only dim=0 is supported.");


inapplicable_msg::InvalidDim?

AnaghaRaoAMD · 2026-01-30T07:37:20Z

-        return false;
+    MIOPEN_SOLVER_INAPPLICABLE_IF(!problem.IsAllContiguous(), inapplicable_msg::NotContiguous);
+
+    MIOPEN_SOLVER_INAPPLICABLE_IF((problem.GetDim() != 0), "Only dim=0 is supported.");


inapplicable_msg::InvalidDim?

AnaghaRaoAMD · 2026-01-30T07:38:29Z

-    if(problem.GetInputDesc().GetNumDims() > 5)
-        return false;
+    MIOPEN_SOLVER_INAPPLICABLE_IF(!IsImprovementOverROCm(problem),
+                                  "Not an improvement over ROCm implementation");


Suggested change

"Not an improvement over ROCm implementation");

inapplicable_msg::IsImprovementOverROCm);

AnaghaRaoAMD · 2026-01-30T07:38:53Z

+    MIOPEN_SOLVER_INAPPLICABLE_IF(!IsImprovementOverROCm(problem),
+                                  "Not an improvement over ROCm implementation");
+
+    MIOPEN_SOLVER_INAPPLICABLE_IF((problem.GetInputDesc().GetNumDims() > 5),


InvalidDim?

AnaghaRaoAMD · 2026-01-30T08:02:10Z

-    if(activ_op.activMode != miopenActivationRELU)
-        return false;
+    MIOPEN_SOLVER_INAPPLICABLE_IF((activ_op.activMode != miopenActivationRELU),
+                                  "Activation is not RELU");


Suggested change

"Activation is not RELU");

inapplicable_msg::InvalidActivation);

AnaghaRaoAMD · 2026-01-30T08:02:53Z

+    MIOPEN_SOLVER_INAPPLICABLE_IF((activationType != miopenActivationRELU &&
+                                   activationType != miopenActivationCLIPPEDRELU &&
+                                   activationType != miopenActivationCLAMP),
+                                  "Unsupported activation type");


Suggested change

"Unsupported activation type");

inapplicable_msg::InvalidActivation);

AnaghaRaoAMD · 2026-01-30T08:03:36Z

+    MIOPEN_SOLVER_INAPPLICABLE_IF((activationType != miopenActivationRELU &&
+                                   activationType != miopenActivationCLIPPEDRELU &&
+                                   activationType != miopenActivationCLAMP),
+                                  "Unsupported activation type");


Suggested change

"Unsupported activation type");

inapplicable_msg::InvalidActivation);

AnaghaRaoAMD · 2026-01-30T08:04:34Z

+        MIOPEN_SOLVER_INAPPLICABLE_IF(!(prim == miopenFusionOpBatchNormInference ||
+                                        prim == miopenFusionOpBiasForward ||
+                                        prim == miopenFusionOpActivForward),
+                                      "Unsupported 2nd fusion op");


Suggested change

"Unsupported 2nd fusion op");

inapplicable_msg::InvalidActivation);

AnaghaRaoAMD · 2026-01-30T08:05:09Z

-        if(!(prim == miopenFusionOpActivForward || prim == miopenFusionOpBatchNormInference))
-            return false;
+        MIOPEN_SOLVER_INAPPLICABLE_IF(
+            !(prim == miopenFusionOpActivForward || prim == miopenFusionOpBatchNormInference),


ThirdFusionOp

AnaghaRaoAMD · 2026-01-30T08:11:58Z

Great work 🙌

As a unit test for logging can we check deterministic msgs being dumped for CPU_UnitTestConvSolverImplicitGemm3DGroupFwdXdlopsDeterministicApplicability_NONE test case?
As a backlog task, can this be leveraged to test applicability for all solvers, so that when it is unintentionally updated we have test coverage?

github-actions · 2026-04-22T21:24:17Z

This pull request has been inactive for 25 days and will be marked as stale.

If you would like to keep this PR open, please:

Add new commits
Add a comment explaining why it should remain open

This PR will be automatically closed in 5 days if no further activity occurs.

github-actions · 2026-04-30T01:06:45Z

This pull request has been automatically closed due to inactivity (30 days with no updates).

If you'd like to continue working on this, feel free to reopen the PR or create a new one.

SreecharanGundaboluAMD added 30 commits November 12, 2025 14:40

implement log and return

5e7d049

conv_ocl_dir2Dfwd.cpp

be09972

Merge remote-tracking branch 'origin/develop' into users/SreecharanGu…

8e131ff

…ndaboluAMD/log_applicability_failures

moved logging to solver_utils

470d8dc

ConvAsm1x1UV2

2795f2f

ConvOclDirectFwd

e29bf4b

implement macros and messages

dd3fac4

activation functions

03484b5

adam

a2a8f90

Batchnorm

0aaf4ec

CatForward

4f8363f

ASM kernels

3d9eacc

Winograd kernels

e7a7cdb

IGEMM ASM

fe0a237

Naive Conv

883fa8c

MLIR

5447019

OCL Solvers

cf86040

Implicit GEMM Solvers

ecadf5c

GEMM solvers

4d092b6

FFT

b9b9f8f

Getitem

67c635a

GLU

17f5d9c

Groupnorm

a52dc98

kth Value

44b4eb7

Layernorm

1f27a43

MHA

ee59806

Multimargin loss

0a44e91

Pooling

d06d8d8

PRELU

f5aeba8

Reduce

18c51d3

SreecharanGundaboluAMD added 6 commits January 20, 2026 01:31

Rope

75852c1

Softmarginloss

b3afeeb

softmax

b3546e1

TensorOP

6243ed4

Fused solvers

3b00290

Merge branch 'develop' into users/SreecharanGundaboluAMD/log_applicab…

7a67baf

…ility_failures

SreecharanGundaboluAMD self-assigned this Jan 20, 2026

SreecharanGundaboluAMD requested a review from a team as a code owner January 20, 2026 02:39

github-actions Bot added the project: miopen label Jan 20, 2026

Update solver_utils.hpp

ac872e0

assistant-librarian Bot added the organization: ROCm label Jan 20, 2026

SreecharanGundaboluAMD added 3 commits January 20, 2026 16:59

clang format fix

accd980

fft solver fix

c2a177a

fix typo

9e904e3

AnaghaRaoAMD reviewed Jan 30, 2026

View reviewed changes

github-actions Bot added the Stale PR has no activity for period of 2+ weeks label Apr 22, 2026

github-actions Bot closed this Apr 30, 2026

	inline constexpr const char* UseAsmKernels = "Assembly kernels are disabled.";
	inline constexpr const char* AsmKernelsDisabled = "Assembly kernels are disabled.";


		MIOPEN_SOLVER_INAPPLICABLE_IF(!ctx.rmv.IsV2orV3(), inapplicable_msg::MetaData);

		MIOPEN_SOLVER_INAPPLICABLE_IF(!(problem.IsFp32() \|\| problem.IsFp16()),


		if(problem.IsTensorsCasted() \|\| problem.IsFp8() \|\| problem.IsBfp8())
		return false;


		MIOPEN_SOLVER_INAPPLICABLE_IF(!problem.IsLayoutDefault(), inapplicable_msg::Layout);

		MIOPEN_SOLVER_INAPPLICABLE_IF(problem.IsTensorsCasted() \|\| problem.IsFp8() \|\| problem.IsBfp8(),

	"Datatype and HW combination not supported");
	"Datatype and device combination not supported");

	"Datatype, HW combination not supported");
	"Datatype and device combination not supported");

	"Legacy CK does not support HW");
	"Legacy CK is not support on this Hardware");

	MIOPEN_SOLVER_INAPPLICABLE_IF(test_cast(problem.GetIn()), "Input tensor missing cast type");
	MIOPEN_SOLVER_INAPPLICABLE_IF(test_cast(problem.GetOut()), "Output tensor missing cast type");

	MIOPEN_SOLVER_INAPPLICABLE_IF(!problem.IsLayoutDefault() && !problem.IsLayoutNHWC(),
	MIOPEN_SOLVER_INAPPLICABLE_IF(!ConvDirectNaiveConvIsApplicableByKernelType(ctx, problem),
	inapplicable_msg::Generic);

		MIOPEN_SOLVER_INAPPLICABLE_IF(!problem.IsLayoutDefault() && !problem.IsLayoutNHWC(),
		inapplicable_msg::Layout);

	"Weight tensor missing cast type");
	"Output tensor missing cast type");

	inapplicable_msg::NoKernelForConfig);
	inapplicable_msg::InvalidGridGemm);

	"GEMM not applicable for F8 on this GPU architecture");
	"GEMM not applicable for FP8 on this GPU architecture");

	MIOPEN_SOLVER_INAPPLICABLE_IF(!IsLargeIndex(problem), "Not a large index problem");
	MIOPEN_SOLVER_INAPPLICABLE_IF(!IsLargeIndex(problem), inapplicable_msg::IndexRange);

	"Not an improvement over ROCm implementation");
	inapplicable_msg::IsImprovementOverROCm);

	"Activation is not RELU");
	inapplicable_msg::InvalidActivation);

	"Unsupported activation type");
	inapplicable_msg::InvalidActivation);

	"Unsupported 2nd fusion op");
	inapplicable_msg::InvalidActivation);

Conversation

SreecharanGundaboluAMD commented Jan 20, 2026

Submission Checklist

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!