refactor: deepen modules and consolidate utilities

- Delete unused CublasSgemm class (duplicate of SGEMMVerifier)
- Move BenchmarkResult to benchmark_core.cuh for kernel layer access
- Move kTensorCoreVerifyTolerance to verify.cuh (centralize tolerances)
- Refactor SGEMMBenchmark to use measureGpuTime (unified timing)
- Refactor tensor_core_benchmark to use measureGpuTime and calculateSgemmMetrics
- Replace int return codes with ParseResult enum (self-documenting API)

Architecture improvements:
- Reduce shallow modules by eliminating duplicate timing logic
- Kernel layer can now use BenchmarkResult without depending on high-level orchestrator
- All verification tolerances now in verify.cuh (improved locality)

Author: LessUp

CONTEXT.md

@@ -57,13 +57,6 @@ Tensor Core 特有的 benchmark 功能，提供：
 - `getTheoreticalPeakGflops()` / `getTheoreticalPeakBandwidth()` - 理论峰值查询
 - `calculateEfficiency()` / `calculateBandwidthUtilization()` - 效率计算
 
-### Benchmark cuBLAS
-**位置**: `src/utils/benchmark_cublas.cuh`
-
-cuBLAS 参考实现：
-- `CublasSgemm` - cuBLAS SGEMM 参考调用器
-- `SgemmReferenceCalculator` - 完整参考计算流程
-
 ### 高级接口
 **位置**: `src/utils/benchmark.cuh`
 

src/cli_parser.cuh

@@ -35,6 +35,21 @@ const std::vector<std::tuple<int, int, int>> BenchmarkConfig::DEFAULT_CASES = {
     {511, 513, 1025},
 };
 
+// ============================================================================
+// CLI 解析结果
+// ============================================================================
+
+/**
+ * CLI 解析结果枚举
+ *
+ * 自文档化的返回类型，替代 int 返回码。
+ */
+enum class ParseResult {
+    Success,    // 解析成功，继续执行
+    Error,      // 解析错误，返回错误码
+    HelpShown   // 显示帮助后正常退出
+};
+
 // ============================================================================
 // CLI 解析器
 // ============================================================================
@@ -86,28 +101,26 @@ class CliParser {
      * 解析命令行参数
      *
      * @param config 输出配置对象
-     * @return 0 成功，1 错误，2 显示帮助后退出
+     * @return ParseResult 枚举指示解析结果
      */
-    int parse(BenchmarkConfig &config) {
+    ParseResult parse(BenchmarkConfig &config) {
         for (int i = 1; i < argc_; ++i) {
             std::string arg = argv_[i];
 
             if (arg == "-h" || arg == "--help") {
                 printUsage(argv_[0]);
-                return 2;
+                return ParseResult::HelpShown;
             }
 
             if (arg == "-s" || arg == "--size") {
-                int result = parseSizeArg(i, config);
-                if (result != 0)
-                    return result;
+                if (!parseSizeArg(i, config))
+                    return ParseResult::Error;
                 continue;
             }
 
             if (arg == "--dims") {
-                int result = parseDimsArg(i, config);
-                if (result != 0)
-                    return result;
+                if (!parseDimsArg(i, config))
+                    return ParseResult::Error;
                 continue;
             }
 
@@ -117,30 +130,28 @@ class CliParser {
             }
 
             if (arg == "--warmup") {
-                int result = parseWarmupArg(i, config);
-                if (result != 0)
-                    return result;
+                if (!parseWarmupArg(i, config))
+                    return ParseResult::Error;
                 continue;
             }
 
             if (arg == "--benchmark") {
-                int result = parseBenchmarkArg(i, config);
-                if (result != 0)
-                    return result;
+                if (!parseBenchmarkArg(i, config))
+                    return ParseResult::Error;
                 continue;
             }
 
             fprintf(stderr, "Unknown argument: %s\n", arg.c_str());
             printUsage(argv_[0]);
-            return 1;
+            return ParseResult::Error;
         }
 
         // 默认添加 1024x1024x1024
         if (config.empty()) {
             config.addCase(1024, 1024, 1024);
         }
 
-        return 0;
+        return ParseResult::Success;
     }
 
     void printUsage(const char *program) const {
@@ -163,82 +174,82 @@ class CliParser {
     }
 
   private:
-    int parseSizeArg(int &i, BenchmarkConfig &config) {
+    bool parseSizeArg(int &i, BenchmarkConfig &config) {
         if (i + 1 >= argc_) {
             fprintf(stderr, "Error: -s requires a size argument\n");
-            return 1;
+            return false;
         }
 
         int size;
         if (!detail::safeStrToInt(argv_[++i], &size, "size")) {
-            return 1;
+            return false;
         }
         if (size <= 0) {
             fprintf(stderr, "Error: Size must be positive\n");
-            return 1;
+            return false;
         }
 
         config.addCase(size, size, size);
-        return 0;
+        return true;
     }
 
-    int parseDimsArg(int &i, BenchmarkConfig &config) {
+    bool parseDimsArg(int &i, BenchmarkConfig &config) {
         if (i + 3 >= argc_) {
             fprintf(stderr, "Error: --dims requires M K N arguments\n");
-            return 1;
+            return false;
         }
 
         int M, K, N;
         if (!detail::safeStrToInt(argv_[++i], &M, "M dimension") ||
             !detail::safeStrToInt(argv_[++i], &K, "K dimension") ||
             !detail::safeStrToInt(argv_[++i], &N, "N dimension")) {
-            return 1;
+            return false;
         }
         if (M <= 0 || K <= 0 || N <= 0) {
             fprintf(stderr, "Error: Dimensions must be positive\n");
-            return 1;
+            return false;
         }
 
         config.addCase(M, K, N);
-        return 0;
+        return true;
     }
 
-    int parseWarmupArg(int &i, BenchmarkConfig &config) {
+    bool parseWarmupArg(int &i, BenchmarkConfig &config) {
         if (i + 1 >= argc_) {
             fprintf(stderr, "Error: --warmup requires a number argument\n");
-            return 1;
+            return false;
         }
 
         int warmup;
         if (!detail::safeStrToInt(argv_[++i], &warmup, "warmup")) {
-            return 1;
+            return false;
         }
         if (warmup < 0) {
             fprintf(stderr, "Error: Warmup runs must be non-negative\n");
-            return 1;
+            return false;
         }
 
         config.warmup_runs = warmup;
-        return 0;
+        return true;
     }
 
-    int parseBenchmarkArg(int &i, BenchmarkConfig &config) {
+    bool parseBenchmarkArg(int &i, BenchmarkConfig &config) {
         if (i + 1 >= argc_) {
             fprintf(stderr, "Error: --benchmark requires a number argument\n");
-            return 1;
+            return false;
         }
 
         int bench;
         if (!detail::safeStrToInt(argv_[++i], &bench, "benchmark")) {
-            return 1;
+            return false;
         }
         if (bench <= 0) {
             fprintf(stderr, "Error: Benchmark runs must be positive\n");
-            return 1;
+            return false;
         }
 
         config.benchmark_runs = bench;
-        return 0;
+        return true;
     }
 
     int argc_;

src/kernels/tensor_core_benchmark.cuh

@@ -1,6 +1,10 @@
 #pragma once
 
 #include "tensor_core_sgemm.cuh"
+#include "../utils/benchmark_core.cuh"
+#include "../utils/benchmark_metrics.cuh"
+#include "../utils/verify.cuh"
+
 #include <cublas_v2.h>
 #include <cuda_fp16.h>
 #include <cuda_runtime.h>
@@ -71,36 +75,17 @@ runTensorCoreComputeOnlyBenchmark(cublasHandle_t cublas_handle, int M, int K, in
     CUDA_CHECK(cudaGetLastError());
     CUDA_CHECK(cudaDeviceSynchronize());
 
-    for (int i = 0; i < warmup_runs; ++i) {
-        d_C.zero();
-        launch_tensor_core_sgemm_fp16(d_A_fp16.get(), d_B_fp16.get(), d_C.get(), M, K, N);
-    }
-    CUDA_CHECK(cudaDeviceSynchronize());
-
-    cudaEvent_t start, stop;
-    CUDA_CHECK(cudaEventCreate(&start));
-    CUDA_CHECK(cudaEventCreate(&stop));
-
-    CUDA_CHECK(cudaEventRecord(start));
-    for (int i = 0; i < benchmark_runs; ++i) {
-        launch_tensor_core_sgemm_fp16(d_A_fp16.get(), d_B_fp16.get(), d_C.get(), M, K, N);
-    }
-    CUDA_CHECK(cudaEventRecord(stop));
-    CUDA_CHECK(cudaEventSynchronize(stop));
-
-    float total_time_ms;
-    CUDA_CHECK(cudaEventElapsedTime(&total_time_ms, start, stop));
-
-    cudaEventDestroy(start);
-    cudaEventDestroy(stop);
-
-    // 填充性能指标
-    result.time_ms = total_time_ms / benchmark_runs;
-    double flops = 2.0 * result.M * result.N * result.K;
-    result.gflops = (flops / (result.time_ms * 1e-3)) / 1e9;
-    double bytes =
-        (result.M * result.K + result.K * result.N + result.M * result.N) * sizeof(float);
-    result.bandwidth_gb_s = (bytes / (result.time_ms * 1e-3)) / 1e9;
+    // 使用统一的 measureGpuTime 计时
+    float time_ms =
+        measureGpuTime([&]() { launch_tensor_core_sgemm_fp16(d_A_fp16.get(), d_B_fp16.get(), d_C.get(), M, K, N); },
+                       warmup_runs, benchmark_runs);
+
+    // 计算指标
+    PerformanceMetrics metrics = calculateSgemmMetrics(M, K, N, time_ms);
+    result.time_ms = metrics.time_ms;
+    result.gflops = metrics.gflops;
+    result.bandwidth_gb_s = metrics.bandwidth_gb_s;
+    result.efficiency = calculateEfficiency(result.gflops, getTheoreticalPeakGflops());
 
     d_C.copyToHost(h_C.data(), M * N);
     d_C_ref.copyToHost(h_C_ref.data(), M * N);

src/kernels/tensor_core_sgemm.cuh

@@ -9,15 +9,18 @@
 // - WMMA FP16→FP32 计算
 // - FP32→FP16 类型转换
 // - 统一启动接口（强制显式 fallback）
-// - 验证容差常量
 //
 // 设计原则：
 // - 深层模块：小接口，大实现
 // - 不提供默认 fallback，强制调用者显式指定
 // - 消除与具体内核的循环依赖
+//
+// 验证容差：
+// - 使用 verify.cuh 中定义的 kTensorCoreVerifyTolerance
 // ============================================================================
 
 #include "../utils/cuda_utils.cuh"
+#include "../utils/verify.cuh"
 #include <cuda_fp16.h>
 #include <cuda_runtime.h>
 #include <functional>
@@ -80,14 +83,6 @@ inline const char *getTensorCoreArchName() {
     return "Unknown";
 }
 
-// ============================================================================
-// Tensor Core Verification Tolerance
-// ============================================================================
-
-// Tensor Core 使用 FP16 中间精度，需要更宽松的容差
-// 此容差定义在 Tensor Core 模块中，保持精度相关常量与其实现在一起
-inline constexpr VerifyTolerance kTensorCoreVerifyTolerance{5e-2f, 1e-2f};
-
 // ============================================================================
 // Fallback 策略接口
 // ============================================================================

src/main.cu

@@ -17,14 +17,14 @@ int main(int argc, char **argv) {
     BenchmarkConfig config;
     CliParser parser(argc, argv);
 
-    int parse_result = parser.parse(config);
-    if (parse_result == 2) {
-        // 显示帮助后正常退出
+    ParseResult result = parser.parse(config);
+    switch (result) {
+    case ParseResult::HelpShown:
         return 0;
-    }
-    if (parse_result != 0) {
-        // 解析错误
+    case ParseResult::Error:
         return 1;
+    case ParseResult::Success:
+        break;
     }
 
     BenchmarkRunner runner(config);