KernelTuningNet for CK [Update CK Fwd, CK Wrw + New integration CK Bwd] #3464

msaudulhassan · 2025-01-08T22:16:05Z

No description provided.

BrianHarrisonAMD · 2025-01-09T15:22:48Z

src/solver/conv/conv_hip_implicit_gemm_grouped_bwd_xdlops.cpp

+        features[0]           = 0.0;
+        features[n + 1]       = problem.GetOutChannels();
+        features[2 * n + 2]   = problem.GetOutHeight();
+        features[3 * n + 3]   = problem.GetOutWidth();
+        features[4 * n + 4]   = problem.GetInChannels();
+        features[5 * n + 5]   = problem.GetInHeight();
+        features[6 * n + 6]   = problem.GetInWidth();
+        features[7 * n + 7]   = problem.GetWeightsHeight();
+        features[8 * n + 8]   = problem.GetWeightsWidth();
+        features[9 * n + 9]   = problem.GetPadH();
+        features[10 * n + 10] = problem.GetPadW();
+        features[11 * n + 11] = problem.GetKernelStrideH();
+        features[12 * n + 12] = problem.GetKernelStrideW();
+        features[13 * n + 13] = problem.GetDilationH();
+        features[14 * n + 14] = problem.GetDilationW();
+        features[15 * n + 15] = problem.GetBatchSize();
+        features[16 * n + 16] = problem.GetInDataType() == miopenFloat ? 2.0 : 1.0;
+        features[17 * n + 17] = problem.GetGroupCount();


For the other arch's, it looks like it's flipping the in/out depending on direction, but not for gfx90a.
Is that intentional?

I suspect that gfx90a should also be flipping in/out based on direction.

BrianHarrisonAMD · 2025-01-09T15:23:51Z

src/solver/conv/conv_hip_implicit_gemm_grouped_bwd_xdlops.cpp

+    //std::cout << "\nBEFORE REMOVING\n";
+    //for (auto i: heuristic_indexes) {
+    //    for (auto x: heuristic_kernels[heuristic_indexes[i]])
+    //        std::cout << x << " ";
+    //    std::cout << "\n";
+    //}


Commented out code could be removed.

BrianHarrisonAMD · 2025-01-09T15:24:16Z

src/solver/conv/conv_hip_implicit_gemm_grouped_bwd_xdlops.cpp

+        std::cout << "\nAFTER REMOVING\n";
+        for (auto i: heuristic_indexes) {
+            for (auto x: heuristic_kernels[heuristic_indexes[i]])
+                std::cout << x << " ";
+            std::cout << "\n";
+            std::cout << valid_kernels[heuristic_indexes[i]];
+            std::cout << "\n";
+        }
+        std::cout << "\n";


Should this be removed, or made into logging statements maybe?

BrianHarrisonAMD · 2025-01-09T15:25:09Z

src/conv/heuristics/ai_heuristics.cpp

+            std::cout << "\n";
+            std::cout << i;
+            std::cout << " " + std::to_string(token) + " " + value;
+            //std::cout << "\nTOKEN SCORES";
+            //for(int j = 0; j < token_scores.size(); j++) {
+            //    std::cout << token_scores[j] << " ";
+            //}
+            //std::cout << std::endl;


Commented out code, and cout.
Should this be removed, or made into logging statements?

BrianHarrisonAMD · 2025-01-09T15:25:33Z

src/conv/heuristics/ai_heuristics.cpp

+    std::cout << "\nFeatures: ";
+    for (int i = 0; i < 17; ++i)
+        std::cout << features[i * 17 + i] << ", ";
+    std::cout << "\n";


Commented out code, and cout.
Should this be removed, or made into logging statements?

BrianHarrisonAMD · 2025-01-09T15:25:46Z

src/conv/heuristics/ai_heuristics.cpp

+    std::cout << "\n PREDICT TYPE: " << model->metadata.predict_type << "\n";
+


Should this be removed, or made into logging statements?

BrianHarrisonAMD · 2025-01-09T15:27:22Z

Looks like the build is failing on formatting check.

Need to run the formatting script:

#!/bin/bash
find . -iname '*.h' \
    -o -iname '*.hpp' \
    -o -iname '*.cpp' \
    -o -iname '*.h.in' \
    -o -iname '*.hpp.in' \
    -o -iname '*.cpp.in' \
    -o -iname '*.cl' \
| grep -v -E '(build/)|(install/)' \

msaudulhassan added 2 commits December 31, 2024 16:47

ck fwd, wrw, bwd KTN updated

0fe2904

solvers merge

201832d

msaudulhassan self-assigned this Jan 8, 2025

msaudulhassan requested review from BrianHarrisonAMD, junliume and BradPepersAMD as code owners January 8, 2025 22:16

BrianHarrisonAMD reviewed Jan 9, 2025

View reviewed changes

BrianHarrisonAMD requested a review from cderb January 9, 2025 15:26

cderb added 2 commits January 9, 2025 11:16

format, convert cout to logs

a3a0cf7

hip tidy

9709e57

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

KernelTuningNet for CK [Update CK Fwd, CK Wrw + New integration CK Bwd] #3464

KernelTuningNet for CK [Update CK Fwd, CK Wrw + New integration CK Bwd] #3464

msaudulhassan commented Jan 8, 2025

BrianHarrisonAMD Jan 9, 2025 •

edited

Loading

cderb Jan 9, 2025 •

edited

Loading

BrianHarrisonAMD Jan 9, 2025

BrianHarrisonAMD Jan 9, 2025

BrianHarrisonAMD Jan 9, 2025

BrianHarrisonAMD Jan 9, 2025

BrianHarrisonAMD Jan 9, 2025 •

edited

Loading

BrianHarrisonAMD commented Jan 9, 2025

		std::cout << "\n PREDICT TYPE: " << model->metadata.predict_type << "\n";

KernelTuningNet for CK [Update CK Fwd, CK Wrw + New integration CK Bwd] #3464

Are you sure you want to change the base?

KernelTuningNet for CK [Update CK Fwd, CK Wrw + New integration CK Bwd] #3464

Conversation

msaudulhassan commented Jan 8, 2025

BrianHarrisonAMD Jan 9, 2025 • edited Loading

Choose a reason for hiding this comment

cderb Jan 9, 2025 • edited Loading

Choose a reason for hiding this comment

BrianHarrisonAMD Jan 9, 2025

Choose a reason for hiding this comment

BrianHarrisonAMD Jan 9, 2025

Choose a reason for hiding this comment

BrianHarrisonAMD Jan 9, 2025

Choose a reason for hiding this comment

BrianHarrisonAMD Jan 9, 2025

Choose a reason for hiding this comment

BrianHarrisonAMD Jan 9, 2025 • edited Loading

Choose a reason for hiding this comment

BrianHarrisonAMD commented Jan 9, 2025

BrianHarrisonAMD Jan 9, 2025 •

edited

Loading

cderb Jan 9, 2025 •

edited

Loading

BrianHarrisonAMD Jan 9, 2025 •

edited

Loading