[MIOpen] Add JSON performance logs for MIOpen convolution driver commands

projects/miopen/driver/README.md

@@ -147,3 +147,89 @@ MIOpenDriver example usages:
 `./bin/MIOpenDriver *base_arg* -?` **OR**  `./bin/MIOpenDriver *base_arg* -h (--help)`
 
 Note: By default the CPU verification is turned on. Verification can be disabled using `-V 0`.
+
+## Environment Variables
+
+### Kernel Name and Execution Time Logging
+
+The `MIOPEN_PERFORMANCE_LOGS` environment variable enables lightweight logging of kernel names and their execution times during MIOpenDriver runs. This is useful for debugging, performance analysis, and understanding which kernels are being executed under different configurations (e.g., different `MIOPEN_FIND_MODE` or `MIOPEN_FORCE` settings). It is meant for usage only with MIOpen driver commands and is untested for direct library usage.
+
+**Logging Levels:**
+
+The variable supports five levels with varying detail and scope:
+
+- **Level 0** (default): No kernel logging
+- **Level 1**: Log only executed solution with **total solution time** - excludes find/search
+- **Level 2**: Log **all kernels** for executed solutions individually - includes transpose/transform kernels, excludes find/search
+- **Level 3**: Log only **performance configs** per solution with **total solution time** - includes find/search
+- **Level 4**: Log **all kernels** individually - includes transpose/transform kernels and find/search kernels
+
+**Usage Examples:**
+
+```bash
+# Level 4: Log all available kernel/performance config/solution information
+export MIOPEN_PERFORMANCE_LOGS=4
+./bin/MIOpenDriver conv -W 32 -H 32 -c 3 -k 32 -x 5 -y 5 -p 2 -q 2
+```
+
+**JSON Output Format:**
+
+Each solution outputs a JSON object with performance configs. Each config contains aggregated timing statistics and optionally individual kernel execution data:
+
+```json
+{
+  "solution": "ConvAsmImplicitGemmGTCDynamicFwdXdlopsNHWC",
+  "solver_id": 1234567890,
+  "phase": "execution",
+  "performance_configs": [
+    {
+      "config_name": "igemm_fwd_gtcx3_nhwc_bf16_bx0_ex1_bt128x128x32...",
+      "config_descriptor": "gemm_m_per_block=128 gemm_n_per_block=128...",
+      "exec_number": 1,
+      "time_executions_ms": [1.52, 1.49, 1.41, 1.38],
+      "time_ms": 1.42,
+      "time_std_ms": 0.056,
+      "time_min_ms": 1.38,
+      "time_max_ms": 1.52,
+      "number_of_transformations": 4,
+      "kernels": [
+        {
+          "kernel_name": "SubTensorOpWithScalar1d",
+          "time_executions_ms": [0.083, 0.104, 0.089],
+          "is_transformation": true
+        },
+        {
+          "kernel_name": "igemm_fwd_gtcx3_nhwc_bf16_bx0_ex1_bt128x128x32...",
+          "time_executions_ms": [0.893, 0.763, 0.724],
+          "is_transformation": false
+        }
+      ]
+    }
+  ]
+}
+```
+
+**JSON Fields:**
+
+**Solution Level:**
+- `solution`: Name of the solver/algorithm
+- `solver_id`: Numeric identifier for the solver
+- `phase`: Either "execution" (actual computation) or "tuning" (find/search phase), "validation", "solver_tuning", or "unknown"
+- `performance_configs`: Array of performance configuration results
+
+**Performance Config Level:**
+- `config_name`: Name of the main kernel or solution
+- `config_descriptor`: Performance config parameters (optional)
+- `exec_number`: Execution counter (starts at 1)
+- `time_executions_ms`: Array of all execution times for this config
+- `time_ms`: Mean execution time (computed using outlier removal)
+- `time_std_ms`: Standard deviation of execution times
+- `time_min_ms`: Minimum execution time
+- `time_max_ms`: Maximum execution time
+- `number_of_transformations`: Count of transformation/transpose kernels
+- `kernels`: Array of individual kernel data (only present in levels 2 & 4, null in levels 1 & 3)
+
+**Kernel Level (when included):**
+- `kernel_name`: Full kernel name
+- `time_executions_ms`: Array of execution times across multiple runs
+- `is_transformation`: Boolean indicating if kernel is a transpose/transform operation

projects/miopen/driver/main.cpp

@@ -1,28 +1,5 @@
-/*******************************************************************************
- *
- * MIT License
- *
- * Copyright (c) 2017 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
- * in the Software without restriction, including without limitation the rights
- * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
- * copies of the Software, and to permit persons to whom the Software is
- * furnished to do so, subject to the following conditions:
- *
- * The above copyright notice and this permission notice shall be included in all
- * copies or substantial portions of the Software.
- *
- * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
- * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
- * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
- * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
- * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
- * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
- * SOFTWARE.
- *
- *******************************************************************************/
+// Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
+// SPDX-License-Identifier: MIT
 #include "driver.hpp"
 #include "registry_driver_maker.hpp"
 
@@ -48,13 +25,25 @@ int main(int argc, char* argv[])
         exit(0); // NOLINT (concurrency-mt-unsafe)
     }
 
+    const bool json_mode = miopen::IsPerformanceLoggingEnabled();
     try
     {
+
         // show command
-        std::cout << "MIOpenDriver";
-        for(int i = 1; i < argc; i++)
-            std::cout << " " << argv[i];
-        std::cout << std::endl;
+        if(!json_mode)
+        {
+            std::cout << "MIOpenDriver";
+            for(int i = 1; i < argc; i++)
+                std::cout << " " << argv[i];
+            std::cout << std::endl;
+        }
+        else
+        {
+            std::cout << "{\"command\":\"MIOpenDriver";
+            for(int i = 1; i < argc; i++)
+                std::cout << " " << argv[i];
+            std::cout << "\"}" << std::endl;
+        }
 
         std::shared_ptr<Driver> drv;
         for(auto f : rdm::GetRegistry())
@@ -65,7 +54,15 @@ int main(int argc, char* argv[])
         }
         if(drv == nullptr)
         {
-            MIOPEN_THROW(miopenStatusBadParm, "Incorrect BaseArg");
+            if(!json_mode)
+            {
+                MIOPEN_THROW(miopenStatusBadParm, "Incorrect BaseArg");
+            }
+            else
+            {
+                std::cout << "{\"error\":\"Incorrect BaseArg\"}" << std::endl;
+                exit(0); // NOLINT (concurrency-mt-unsafe)
+            }
         }
 
         drv->name = base_arg;
@@ -74,14 +71,33 @@ int main(int argc, char* argv[])
         int rc = drv->ParseCmdLineArgs(argc, argv);
         if(rc != 0)
         {
-            std::cout << "ParseCmdLineArgs() FAILED, rc = " << rc << std::endl;
+            if(!json_mode)
+            {
+                std::cout << "ParseCmdLineArgs() FAILED, rc = " << rc << std::endl;
+            }
+            else
+            {
+                std::cout << "{\"error\":\"ParseCmdLineArgs() FAILED, rc = " << rc << "\"}"
+                          << std::endl;
+            }
             return rc;
         }
+
         drv->GetandSetData();
+
         rc = drv->AllocateBuffersAndCopy();
+
         if(rc != 0)
         {
-            std::cout << "AllocateBuffersAndCopy() FAILED, rc = " << rc << std::endl;
+            if(!json_mode)
+            {
+                std::cout << "AllocateBuffersAndCopy() FAILED, rc = " << rc << std::endl;
+            }
+            else
+            {
+                std::cout << "{\"error\":\"AllocateBuffersAndCopy() FAILED, rc = " << rc << "\"}"
+                          << std::endl;
+            }
             return rc;
         }
 
@@ -108,8 +124,18 @@ int main(int argc, char* argv[])
             rc = drv->RunForwardGPU();
             cumulative_rc |= rc;
             if(rc != 0)
-                std::cout << "RunForwardGPU() FAILED, rc = " << "0x" << std::hex << rc << std::dec
-                          << std::endl;
+            {
+                if(!json_mode)
+                {
+                    std::cout << "RunForwardGPU() FAILED, rc = " << "0x" << std::hex << rc
+                              << std::dec << std::endl;
+                }
+                else
+                {
+                    std::cout << "{\"error\":\"RunForwardGPU() FAILED, rc = 0x" << std::hex << rc
+                              << std::dec << "\"}" << std::endl;
+                }
+            }
             if(verifyarg) // Verify even if Run() failed.
                 cumulative_rc |= drv->VerifyForward();
         }
@@ -119,22 +145,55 @@ int main(int argc, char* argv[])
             rc = drv->RunBackwardGPU();
             cumulative_rc |= rc;
             if(rc != 0)
-                std::cout << "RunBackwardGPU() FAILED, rc = " << "0x" << std::hex << rc << std::dec
-                          << std::endl;
+            {
+                if(!json_mode)
+                {
+                    std::cout << "RunBackwardGPU() FAILED, rc = " << "0x" << std::hex << rc
+                              << std::dec << std::endl;
+                }
+                else
+                {
+                    std::cout << "{\"error\":\"RunBackwardGPU() FAILED, rc = 0x" << std::hex << rc
+                              << std::dec << "\"}" << std::endl;
+                }
+            }
             if(verifyarg) // Verify even if Run() failed.
                 cumulative_rc |= drv->VerifyBackward();
         }
 
+        // Flush any accumulated performance logging data before exiting
+        miopen::FinalizeJsonLogging();
+
         return cumulative_rc;
     }
     catch(const miopen::Exception& ex)
     {
-        std::cerr << "Error: " << ex.what() << std::endl;
+        // Flush any accumulated performance logging data before exiting
+        miopen::FinalizeJsonLogging();
+
+        if(!json_mode)
+        {
+            std::cerr << "Error: " << ex.what() << std::endl;
+        }
+        else
+        {
+            std::cout << "{\"error\":\"" << ex.what() << "\"}" << std::endl;
+        }
         return EXIT_FAILURE;
     }
     catch(const std::exception& ex)
     {
-        std::cerr << "Error: " << ex.what() << std::endl;
+        // Flush any accumulated performance logging data before exiting
+        miopen::FinalizeJsonLogging();
+
+        if(!json_mode)
+        {
+            std::cerr << "Error: " << ex.what() << std::endl;
+        }
+        else
+        {
+            std::cout << "{\"error\":\"" << ex.what() << "\"}" << std::endl;
+        }
         return EXIT_FAILURE;
     }
 }

projects/miopen/driver/random.hpp

@@ -1,32 +1,10 @@
-/*******************************************************************************
- *
- * MIT License
- *
- * Copyright (c) 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
- * in the Software without restriction, including without limitation the rights
- * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
- * copies of the Software, and to permit persons to whom the Software is
- * furnished to do so, subject to the following conditions:
- *
- * The above copyright notice and this permission notice shall be included in all
- * copies or substantial portions of the Software.
- *
- * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
- * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
- * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
- * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
- * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
- * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
- * SOFTWARE.
- *
- *******************************************************************************/
+// Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
+// SPDX-License-Identifier: MIT
 #ifndef GUARD_RANDOM_GEN_
 #define GUARD_RANDOM_GEN_
 
 #include <miopen/env.hpp>
+#include <miopen/kernel_tuning_mode.hpp>
 
 #include <cassert>
 #include <iostream>
@@ -48,7 +26,16 @@ inline std::random_device::result_type get_default_seed()
         auto seed = external_seed == 0
                         ? std::random_device{}()
                         : static_cast<std::random_device::result_type>(external_seed);
-        std::cout << "PRNG seed: " << seed << "\n";
+
+        // Check if JSON mode is enabled in MIOPEN_PERFORMANCE_LOGS
+        if(miopen::IsPerformanceLoggingEnabled())
+        {
+            std::cout << "{\"prng_seed\":" << seed << "}" << std::endl;
+        }
+        else
+        {
+            std::cout << "PRNG seed: " << seed << "\n";
+        }
         return seed;
     }()};
     return seed;