Skip to content

Commit c5c7d0d

Browse files
Adding stress test for autogeneration and autocorrection, and fixing preshuffle bug.
1 parent b386bfa commit c5c7d0d

33 files changed

+3444
-413
lines changed

dispatcher/examples/conv/cpp/02_conv_validation.cpp

Lines changed: 30 additions & 20 deletions
Original file line numberDiff line numberDiff line change
@@ -2,13 +2,13 @@
22
// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved.
33

44
/**
5-
* Example 03: Convolution with CPU Validation - Declarative
5+
* Example 02: Convolution with CPU Validation - Declarative
66
*
77
* Demonstrates convolution with CPU reference verification.
88
* Uses the Signature/Algorithm/Arch declarative pattern.
99
*
1010
* Self-contained build:
11-
* python3 scripts/compile_conv_examples.py examples/conv/cpp/03_conv_validation.cpp
11+
* python3 scripts/compile_conv_examples.py examples/conv/cpp/02_conv_validation.cpp
1212
*
1313
* Complexity: ★★★☆☆
1414
*/
@@ -22,6 +22,7 @@
2222

2323
// Declarative utilities
2424
#include "ck_tile/dispatcher/conv_utils.hpp"
25+
#include "ck_tile/dispatcher/example_args.hpp"
2526

2627
// CK Tile includes
2728
#include "ck_tile/core.hpp"
@@ -32,6 +33,7 @@
3233

3334
using namespace ck_tile::dispatcher;
3435
using namespace ck_tile::dispatcher::conv_utils;
36+
using namespace ck_tile::dispatcher::utils;
3537

3638
// =============================================================================
3739
// KERNEL DECLARATIONS
@@ -63,10 +65,28 @@ using AccDataType = float;
6365

6466
int main(int argc, char* argv[])
6567
{
68+
ExampleArgs args("Example 02: Conv Validation", "Convolution with CPU reference verification");
69+
args.add_option("-n", "1", "Batch size N");
70+
args.add_option("-c", "64", "Input channels C");
71+
args.add_option("-k", "128", "Output channels K");
72+
args.add_option("--size", "14", "Spatial size (H=W)");
73+
args.add_flag("--no-verify", "Skip CPU validation");
74+
args.add_flag("--list", "List all kernel sets");
75+
76+
if(!args.parse(argc, argv))
77+
return 0;
78+
6679
std::cout << "======================================================================\n";
67-
std::cout << "Example 03: Convolution with CPU Validation (Declarative)\n";
80+
std::cout << "Example 02: Convolution with CPU Validation (Declarative)\n";
6881
std::cout << "======================================================================\n\n";
6982

83+
if(args.has("--list"))
84+
{
85+
std::cout << "Declared Kernel Sets:\n";
86+
ConvKernelSetRegistry::instance().print();
87+
return 0;
88+
}
89+
7090
// -------------------------------------------------------------------------
7191
// Step 1: Show declared kernels
7292
// -------------------------------------------------------------------------
@@ -83,23 +103,13 @@ int main(int argc, char* argv[])
83103
std::cout << "Step 2: Define Problem\n";
84104
std::cout << "----------------------\n";
85105

86-
int N = 1, C = 64, K = 128, Hi = 14, Wi = 14, Y = 3, X = 3;
87-
bool verify = true;
88-
89-
for(int i = 1; i < argc; ++i)
90-
{
91-
std::string arg = argv[i];
92-
if(arg == "-n" && i + 1 < argc)
93-
N = std::stoi(argv[++i]);
94-
else if(arg == "-c" && i + 1 < argc)
95-
C = std::stoi(argv[++i]);
96-
else if(arg == "-k" && i + 1 < argc)
97-
K = std::stoi(argv[++i]);
98-
else if(arg == "-h" && i + 1 < argc)
99-
Hi = Wi = std::stoi(argv[++i]);
100-
else if(arg == "--no-verify")
101-
verify = false;
102-
}
106+
int N = args.get_int("-n", 1);
107+
int C = args.get_int("-c", 64);
108+
int K = args.get_int("-k", 128);
109+
int Hi = args.get_int("--size", 14);
110+
int Wi = Hi;
111+
int Y = 3, X = 3;
112+
bool verify = !args.has("--no-verify");
103113

104114
auto problem = create_conv2d_problem(N, C, K, Hi, Wi, Y, X, 1, 1, ConvOp::Forward);
105115
print_problem(problem);

dispatcher/examples/conv/cpp/03_multi_size.cpp

Lines changed: 19 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -2,7 +2,7 @@
22
// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved.
33

44
/**
5-
* Example 04: Multi-Size Convolution with GPU Execution
5+
* Example 03: Multi-Size Convolution with GPU Execution
66
*
77
* Demonstrates using different kernel tile sizes for different problem sizes,
88
* with actual GPU execution for each.
@@ -16,13 +16,15 @@
1616
#include <hip/hip_runtime.h>
1717

1818
#include "ck_tile/dispatcher/conv_utils.hpp"
19+
#include "ck_tile/dispatcher/example_args.hpp"
1920
#include "ck_tile/core.hpp"
2021
#include "ck_tile/host.hpp"
2122
#include "ck_tile/host/convolution_parameter.hpp"
2223
#include "ck_tile/ops/grouped_convolution.hpp"
2324

2425
using namespace ck_tile::dispatcher;
2526
using namespace ck_tile::dispatcher::conv_utils;
27+
using namespace ck_tile::dispatcher::utils;
2628

2729
// =============================================================================
2830
// KERNEL DECLARATIONS - Multiple tile sizes
@@ -131,12 +133,26 @@ void run_conv_on_gpu(const ConvProblem& problem, const std::string& label)
131133
// MAIN
132134
// =============================================================================
133135

134-
int main()
136+
int main(int argc, char* argv[])
135137
{
138+
ExampleArgs args("Example 03: Multi-Size Conv",
139+
"Different tile sizes for different problem sizes");
140+
args.add_flag("--list", "List all kernel sets");
141+
142+
if(!args.parse(argc, argv))
143+
return 0;
144+
136145
std::cout << "======================================================================\n";
137-
std::cout << "Example 04: Multi-Size Convolution with GPU Execution\n";
146+
std::cout << "Example 03: Multi-Size Convolution with GPU Execution\n";
138147
std::cout << "======================================================================\n\n";
139148

149+
if(args.has("--list"))
150+
{
151+
std::cout << "Declared Kernel Sets:\n";
152+
ConvKernelSetRegistry::instance().print();
153+
return 0;
154+
}
155+
140156
// -------------------------------------------------------------------------
141157
// Step 1: Show declared kernels
142158
// -------------------------------------------------------------------------

dispatcher/examples/conv/cpp/05_heuristics.cpp

Lines changed: 18 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -2,7 +2,7 @@
22
// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved.
33

44
/**
5-
* Example 06: Convolution Heuristics with GPU Execution
5+
* Example 05: Convolution Heuristics with GPU Execution
66
*
77
* Demonstrates heuristic-based kernel selection with GPU execution.
88
*
@@ -14,13 +14,15 @@
1414
#include <hip/hip_runtime.h>
1515

1616
#include "ck_tile/dispatcher/conv_utils.hpp"
17+
#include "ck_tile/dispatcher/example_args.hpp"
1718
#include "ck_tile/core.hpp"
1819
#include "ck_tile/host.hpp"
1920
#include "ck_tile/host/convolution_parameter.hpp"
2021
#include "ck_tile/ops/grouped_convolution.hpp"
2122

2223
using namespace ck_tile::dispatcher;
2324
using namespace ck_tile::dispatcher::conv_utils;
25+
using namespace ck_tile::dispatcher::utils;
2426

2527
// =============================================================================
2628
// KERNEL DECLARATIONS
@@ -82,12 +84,25 @@ using OutDataType = ck_tile::half_t;
8284
// MAIN
8385
// =============================================================================
8486

85-
int main()
87+
int main(int argc, char* argv[])
8688
{
89+
ExampleArgs args("Example 05: Conv Heuristics", "Heuristic-based kernel selection");
90+
args.add_flag("--list", "List all kernel sets");
91+
92+
if(!args.parse(argc, argv))
93+
return 0;
94+
8795
std::cout << "======================================================================\n";
88-
std::cout << "Example 06: Convolution Heuristics with GPU Execution\n";
96+
std::cout << "Example 05: Convolution Heuristics with GPU Execution\n";
8997
std::cout << "======================================================================\n\n";
9098

99+
if(args.has("--list"))
100+
{
101+
std::cout << "Declared Kernel Sets:\n";
102+
ConvKernelSetRegistry::instance().print();
103+
return 0;
104+
}
105+
91106
// -------------------------------------------------------------------------
92107
// Setup
93108
// -------------------------------------------------------------------------

dispatcher/examples/conv/cpp/06_json_export.cpp

Lines changed: 22 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -2,7 +2,7 @@
22
// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved.
33

44
/**
5-
* Example 07: Convolution JSON Export with GPU Execution
5+
* Example 06: Convolution JSON Export with GPU Execution
66
*
77
* Exports kernel configurations to JSON and runs on GPU.
88
*
@@ -16,13 +16,15 @@
1616
#include <hip/hip_runtime.h>
1717

1818
#include "ck_tile/dispatcher/conv_utils.hpp"
19+
#include "ck_tile/dispatcher/example_args.hpp"
1920
#include "ck_tile/core.hpp"
2021
#include "ck_tile/host.hpp"
2122
#include "ck_tile/host/convolution_parameter.hpp"
2223
#include "ck_tile/ops/grouped_convolution.hpp"
2324

2425
using namespace ck_tile::dispatcher;
2526
using namespace ck_tile::dispatcher::conv_utils;
27+
using namespace ck_tile::dispatcher::utils;
2628

2729
// =============================================================================
2830
// KERNEL DECLARATIONS
@@ -99,12 +101,26 @@ using OutDataType = ck_tile::half_t;
99101
// MAIN
100102
// =============================================================================
101103

102-
int main()
104+
int main(int argc, char* argv[])
103105
{
106+
ExampleArgs args("Example 06: Conv JSON Export", "Export kernel configurations to JSON");
107+
args.add_option("--output", "conv_kernels.json", "Output JSON file path");
108+
args.add_flag("--list", "List all kernel sets");
109+
110+
if(!args.parse(argc, argv))
111+
return 0;
112+
104113
std::cout << "======================================================================\n";
105-
std::cout << "Example 07: Convolution JSON Export with GPU Execution\n";
114+
std::cout << "Example 06: Convolution JSON Export with GPU Execution\n";
106115
std::cout << "======================================================================\n\n";
107116

117+
if(args.has("--list"))
118+
{
119+
std::cout << "Declared Kernel Sets:\n";
120+
ConvKernelSetRegistry::instance().print();
121+
return 0;
122+
}
123+
108124
// -------------------------------------------------------------------------
109125
// Export to JSON
110126
// -------------------------------------------------------------------------
@@ -117,12 +133,13 @@ int main()
117133
std::cout << json << "\n";
118134

119135
// Write to file
120-
std::ofstream file("conv_kernels.json");
136+
std::string output_file = args.get("--output");
137+
std::ofstream file(output_file);
121138
if(file)
122139
{
123140
file << json;
124141
file.close();
125-
std::cout << "[Saved to conv_kernels.json]\n\n";
142+
std::cout << "[Saved to " << output_file << "]\n\n";
126143
}
127144

128145
// -------------------------------------------------------------------------

dispatcher/examples/conv/cpp/07_multi_registry.cpp

Lines changed: 19 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -2,7 +2,7 @@
22
// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved.
33

44
/**
5-
* Example 08: Multiple Convolution Registries with GPU Execution
5+
* Example 07: Multiple Convolution Registries with GPU Execution
66
*
77
* Demonstrates using separate registries for different use cases,
88
* each running on GPU.
@@ -15,13 +15,15 @@
1515
#include <hip/hip_runtime.h>
1616

1717
#include "ck_tile/dispatcher/conv_utils.hpp"
18+
#include "ck_tile/dispatcher/example_args.hpp"
1819
#include "ck_tile/core.hpp"
1920
#include "ck_tile/host.hpp"
2021
#include "ck_tile/host/convolution_parameter.hpp"
2122
#include "ck_tile/ops/grouped_convolution.hpp"
2223

2324
using namespace ck_tile::dispatcher;
2425
using namespace ck_tile::dispatcher::conv_utils;
26+
using namespace ck_tile::dispatcher::utils;
2527

2628
// =============================================================================
2729
// KERNEL DECLARATIONS - Different registries for different use cases
@@ -119,12 +121,26 @@ float run_conv(int N, int C, int K, int H, int W)
119121
// MAIN
120122
// =============================================================================
121123

122-
int main()
124+
int main(int argc, char* argv[])
123125
{
126+
ExampleArgs args("Example 07: Multi-Registry Conv",
127+
"Separate registries for different use cases");
128+
args.add_flag("--list", "List all kernel sets");
129+
130+
if(!args.parse(argc, argv))
131+
return 0;
132+
124133
std::cout << "======================================================================\n";
125-
std::cout << "Example 08: Multiple Convolution Registries with GPU Execution\n";
134+
std::cout << "Example 07: Multiple Convolution Registries with GPU Execution\n";
126135
std::cout << "======================================================================\n\n";
127136

137+
if(args.has("--list"))
138+
{
139+
std::cout << "Declared Kernel Sets:\n";
140+
ConvKernelSetRegistry::instance().print();
141+
return 0;
142+
}
143+
128144
// -------------------------------------------------------------------------
129145
// Create separate registries
130146
// -------------------------------------------------------------------------

dispatcher/examples/conv/cpp/08_conv3d_forward.cpp

Lines changed: 30 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -2,7 +2,7 @@
22
// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved.
33

44
/**
5-
* Example 09: 3D Convolution Forward with GPU Execution
5+
* Example 08: 3D Convolution Forward with GPU Execution
66
*
77
* Demonstrates 3D convolution (e.g., for video or volumetric data).
88
*
@@ -14,13 +14,15 @@
1414
#include <hip/hip_runtime.h>
1515

1616
#include "ck_tile/dispatcher/conv_utils.hpp"
17+
#include "ck_tile/dispatcher/example_args.hpp"
1718
#include "ck_tile/core.hpp"
1819
#include "ck_tile/host.hpp"
1920
#include "ck_tile/host/convolution_parameter.hpp"
2021
#include "ck_tile/ops/grouped_convolution.hpp"
2122

2223
using namespace ck_tile::dispatcher;
2324
using namespace ck_tile::dispatcher::conv_utils;
25+
using namespace ck_tile::dispatcher::utils;
2426

2527
// =============================================================================
2628
// KERNEL DECLARATIONS - 3D Forward
@@ -56,12 +58,30 @@ using OutDataType = ck_tile::half_t;
5658
// MAIN
5759
// =============================================================================
5860

59-
int main()
61+
int main(int argc, char* argv[])
6062
{
63+
ExampleArgs args("Example 08: Conv3D Forward", "3D convolution for video/volumetric data");
64+
args.add_option("-n", "1", "Batch size N");
65+
args.add_option("-c", "32", "Input channels C");
66+
args.add_option("-k", "64", "Output channels K");
67+
args.add_option("--depth", "8", "Depth D");
68+
args.add_option("--size", "16", "Spatial size (H=W)");
69+
args.add_flag("--list", "List all kernel sets");
70+
71+
if(!args.parse(argc, argv))
72+
return 0;
73+
6174
std::cout << "======================================================================\n";
62-
std::cout << "Example 09: 3D Convolution Forward with GPU Execution\n";
75+
std::cout << "Example 08: 3D Convolution Forward with GPU Execution\n";
6376
std::cout << "======================================================================\n\n";
6477

78+
if(args.has("--list"))
79+
{
80+
std::cout << "Declared Kernel Sets:\n";
81+
ConvKernelSetRegistry::instance().print();
82+
return 0;
83+
}
84+
6585
// -------------------------------------------------------------------------
6686
// Step 1: Show declared kernels
6787
// -------------------------------------------------------------------------
@@ -78,9 +98,13 @@ int main()
7898
std::cout << "Step 2: Define 3D Problem\n";
7999
std::cout << "-------------------------\n";
80100

81-
// 3D problem: N=1, C=32, K=64, D=8, H=16, W=16, filter 3x3x3
82-
int N = 1, C = 32, K = 64;
83-
int Di = 8, Hi = 16, Wi = 16;
101+
// 3D problem from args
102+
int N = args.get_int("-n", 1);
103+
int C = args.get_int("-c", 32);
104+
int K = args.get_int("-k", 64);
105+
int Di = args.get_int("--depth", 8);
106+
int Hi = args.get_int("--size", 16);
107+
int Wi = Hi;
84108
int Z = 3, Y = 3, X = 3;
85109

86110
auto problem = create_conv3d_problem(N, C, K, Di, Hi, Wi, Z, Y, X, 1, 1, ConvOp::Forward);

0 commit comments

Comments
 (0)