TheRock/rocm-libraries/dnn-providers/miopen-provider/docs/Knobs.md

MIOpen Provider - Configuration Knobs

This document describes the configuration knobs available for the MIOpen Provider Plugin for hipDNN.

For general information about hipDNN's knobs system, see the hipDNN Knobs Documentation.

Table of Contents

• Overview
• Available Knobs
• Knob Details
  • Benchmarking
  • Workspace Size Limit
• Usage Examples
• Performance Considerations
• Best Practices

---

Overview

The MIOpen Provider Plugin supports configuration knobs that control kernel selection, performance tuning, and memory usage. These knobs allow you to optimize MIOpen's behavior for your specific workload and hardware configuration.

Knob Types

The MIOpen Provider supports two types of knobs:

• Global Knobs: Standard knobs available for all engines (namespace: global.*)
• Custom Knobs: Operation-specific knobs provided dynamically based on the graph (no custom knobs currently)

---

Available Knobs

The following table lists all configuration knobs supported by the MIOpen Provider:

Knob Name	Type	Scope	Default	Valid Range	Description
global.benchmarking	int64	All Operations	0 (disabled)	0-1	Enable benchmarking mode for kernel selection
global.workspace_size_limit	int64	Convolution Operations Only	Maximum	Dynamic (solver-dependent)	Maximum workspace memory in bytes

Note

The global.workspace_size_limit knob is only available for convolution operations (Forward, Backward Data, Backward Weights). It is not supported for batchnorm or other operations.

---

Knob Details

Benchmarking

Knob Name: global.benchmarking

Purpose: Controls whether MIOpen performs kernel benchmarking to find the optimal solver for a given operation.

Type: Integer (int64)

Valid Values:

• 0 - Benchmarking disabled (default)
• 1 - Benchmarking enabled

Behavior:

When disabled (default):

• MIOpen uses heuristics to select a kernel
• First execution is fast
• May not use the optimal kernel for your specific configuration

When enabled:

• MIOpen benchmarks multiple solver candidates
• First execution is significantly slower (seconds to minutes)
• Subsequent executions use the cached optimal solver
• Provides best performance for production workloads

Caching:

• Benchmark results are cached in MIOpen's performance database
• Default location: ~/.config/miopen/ on Linux
• Cache persists across application runs
• Cache is specific to: GPU model, operation parameters, tensor dimensions, and data types

Performance Impact:

• First run with benchmarking: Can take seconds to minutes depending on operation complexity
• Subsequent runs: Minimal overhead (cache lookup)
• Production benefit: Typically 10-50% performance improvement over heuristic selection

Example:

// Enable benchmarking
std::vector<KnobSetting> settings;
settings.emplace_back("global.benchmarking", 1);
graph.create_execution_plan_ext(engineId, settings);

Workspace Size Limit

Knob Name: global.workspace_size_limit

Purpose: Limits the maximum workspace memory that MIOpen solvers can use for convolution operations.

Type: Integer (int64) - workspace size in bytes

Scope: Convolution operations only (Forward, Backward Data, Backward Weights)

Valid Range: Dynamic, determined at runtime based on available MIOpen solvers

The valid range is operation-specific and depends on:

• Convolution type (Forward/Backward Data/Backward Weights)
• Tensor dimensions and data types
• Available MIOpen solvers for the configuration
• GPU memory constraints

How the Range is Determined:

1. MIOpen queries all available solvers for the specific operation and tensor configuration
2. Each solver reports its workspace memory requirement
3. The minimum workspace is the smallest requirement across all solvers
4. The maximum workspace is the largest requirement across all solvers
5. Default is set to the maximum for optimal performance

Example Range: For a specific convolution forward operation:

• Minimum: 512 KB (lightweight kernel with minimal workspace)
• Maximum: 128 MB (high-performance kernel with large workspace)
• Default: 128 MB (use maximum for best performance)

Behavior:

Setting this knob to a value lower than the maximum:

• Constrains solver selection to only those requiring ≤ the specified workspace
• May reduce performance if optimal solvers require more workspace
• Useful for memory-constrained systems where total GPU memory is limited

Setting this knob to the maximum (or not setting it):

• Allows all solvers to be considered
• Provides best performance
• Uses more GPU memory

Important Notes:

Important

The global.workspace_size_limit knob is dynamically provided only when applicable. It will not appear in the knobs list for non-convolution operations (e.g., batchnorm, pointwise).

Warning

Setting a workspace limit below the minimum required by all solvers will result in an error when creating the execution plan.

Example:

// Query knobs to find valid range
std::vector<Knob> knobs;
graph.get_knobs_for_engine(engineId, knobs);

for (const auto& knob : knobs) {
    if (knob.knobId() == "global.workspace_size_limit") {
        // Check constraints to see valid range
        const auto* constraint = knob.constraint();
        std::cout << "Workspace range: " << constraint->toString() << "\n";
    }
}

// Limit workspace to 32 MB
std::vector<KnobSetting> settings;
settings.emplace_back("global.workspace_size_limit", 32LL * 1024 * 1024);
graph.create_execution_plan_ext(engineId, settings);

---

Usage Examples

Example 1: Query Available Knobs

#include <hipdnn_frontend.hpp>

using namespace hipdnn_frontend;

// After building the graph
Graph graph;
// ... setup and build graph ...

// Query knobs for MIOpen engine
std::vector<Knob> knobs;
auto error = graph.get_knobs_for_engine(MIOPEN_ENGINE_ID, knobs);

if (error.is_good()) {
    std::cout << "Available knobs for MIOpen engine:\n";
    for (const auto& knob : knobs) {
        std::cout << "  " << knob.knobId() << ": " << knob.description() << "\n";
        std::cout << "  Default: ";

        const auto& defaultVal = knob.defaultValue();
        if (std::holds_alternative<int64_t>(defaultVal)) {
            std::cout << std::get<int64_t>(defaultVal);
        }
        std::cout << "\n";

        if (const auto* constraint = knob.constraint()) {
            std::cout << "  Constraint: " << constraint->toString() << "\n";
        }
        std::cout << "\n";
    }
}

Example 2: Combined Knob Settings

#include <hipdnn_frontend.hpp>

using namespace hipdnn_frontend;

Graph graph;
// ... setup convolution graph ...
graph.build_operation_graph(handle);

// Enable benchmarking and set workspace limit
std::vector<KnobSetting> settings;
settings.emplace_back("global.benchmarking", 1);
settings.emplace_back("global.workspace_size_limit", 128LL * 1024 * 1024);

auto error = graph.create_execution_plan_ext(MIOPEN_ENGINE_ID, settings);

if (error.is_good()) {
    std::cout << "Execution plan created with custom knob settings\n";
}

---

Best Practices

For Development

1. Start with defaults: Use default knob values during initial development
2. Profile first: Measure baseline performance before tuning knobs
3. Query knobs: Always check available knobs and their constraints using get_knobs_for_engine()
4. Test incremental changes: Modify one knob at a time to understand impact

For Production

1. Enable benchmarking during warm-up:

   // Warm-up phase
   std::vector<KnobSetting> warmupSettings;
   warmupSettings.emplace_back("global.benchmarking", 1);
   graph.create_execution_plan_ext(engineId, warmupSettings);
   
   // Execute a few times to populate cache
   for (int i = 0; i < 5; i++) {
       graph.execute(handle, variantPack, workspace);
   }
   

2. Use cached results in production: After warm-up, benchmarking can be disabled as results are cached

3. Document knob settings: Keep a record of knob configurations used in production for reproducibility

For Memory-Constrained Environments

1. Query workspace ranges:

   // Find minimum and maximum workspace for the operation
   std::vector<Knob> knobs;
   graph.get_knobs_for_engine(engineId, knobs);
   
   for (const auto& knob : knobs) {
       if (knob.knobId() == "global.workspace_size_limit") {
           // Log constraint to understand valid range
       }
   }
   

2. Set conservative limits: Start with a lower workspace limit and increase if performance is insufficient

3. Balance batch size and workspace: Reducing workspace allows larger batch sizes, which may offset performance loss

Error Handling

Always check for errors when setting knobs:

auto error = graph.create_execution_plan_ext(engineId, settings);
if (!error.is_good()) {
    std::cerr << "Failed to create execution plan: " << error.get_message() << "\n";

    // Common errors:
    // - Workspace limit below minimum
    // - Invalid knob ID
    // - Value outside valid range
}

---

Related Documentation

• hipDNN Knobs Documentation - Comprehensive knobs system guide
• MIOpen Provider Operation Support - Supported operations
• MIOpen Documentation - MIOpen library documentation
• Plugin Development Guide - How to implement knobs in plugins