refactor: improve code safety and quality

- Fix DeviceMemory constructor exception safety (initialize ptr_ to nullptr)
- Add gridSize overflow check for large matrix support
- Extract magic numbers into named constants (kDefaultTileSize, kDefaultBlockSize, kFilenameBufferSize)
- Add DeviceInfoCache singleton to cache device properties and reduce redundant queries
- Add [[maybe_unused]] attribute to utility functions (nullFallback, calculateBandwidthUtilization, printPerformanceReport)

Author: LessUp

src/benchmark_runner.cuh

@@ -103,7 +103,7 @@ class BenchmarkRunner {
         printPerformanceComparison(benchmark.getResults(), cublas_gflops);
 
         // 导出 roofline 数据
-        char filename[256];
+        char filename[kFilenameBufferSize];
         snprintf(filename, sizeof(filename), "roofline_data_%d_%d_%d.csv", M, K, N);
         benchmark.exportRooflineData(filename);
     }

src/kernels/tensor_core_benchmark.cuh

@@ -5,6 +5,7 @@
 #include "../utils/benchmark_metrics.cuh"
 #include "../utils/verify.cuh"
 
+#include <climits>
 #include <cublas_v2.h>
 #include <cuda_fp16.h>
 #include <cuda_runtime.h>
@@ -66,9 +67,17 @@ runTensorCoreComputeOnlyBenchmark(cublasHandle_t cublas_handle, int M, int K, in
     CUBLAS_CHECK(cublasSgemm(cublas_handle, CUBLAS_OP_N, CUBLAS_OP_N, N, M, K, &alpha,
                              d_B.get(), N, d_A.get(), K, &beta, d_C_ref.get(), N));
 
-    int blockSize = 256;
-    int gridSizeA = (M * K + blockSize - 1) / blockSize;
-    int gridSizeB = (K * N + blockSize - 1) / blockSize;
+    int blockSize = kDefaultBlockSize;
+    // 安全计算 gridSize，检查溢出
+    auto safeGridSize = [](size_t num, int blk) -> int {
+        size_t grid = (num + blk - 1) / blk;
+        if (grid > static_cast<size_t>(INT_MAX)) {
+            throw CudaError("Grid size overflow: matrix too large for kernel launch");
+        }
+        return static_cast<int>(grid);
+    };
+    int gridSizeA = safeGridSize(static_cast<size_t>(M) * K, blockSize);
+    int gridSizeB = safeGridSize(static_cast<size_t>(K) * N, blockSize);
 
     float_to_half_kernel<<<gridSizeA, blockSize>>>(d_A.get(), d_A_fp16.get(), M * K);
     float_to_half_kernel<<<gridSizeB, blockSize>>>(d_B.get(), d_B_fp16.get(), K * N);

src/kernels/tensor_core_sgemm.cuh

@@ -21,6 +21,7 @@
 
 #include "../utils/cuda_utils.cuh"
 #include "../utils/verify.cuh"
+#include <climits>
 #include <cuda_fp16.h>
 #include <cuda_runtime.h>
 #include <functional>
@@ -46,14 +47,7 @@ using tensor_core::WMMA_N;
 /**
  * 检查当前设备是否支持 Tensor Core (sm_70+)
  */
-inline bool tensorCoresAvailable() {
-    int device;
-    CUDA_CHECK(cudaGetDevice(&device));
-
-    cudaDeviceProp prop;
-    CUDA_CHECK(cudaGetDeviceProperties(&prop, device));
-    return prop.major >= 7;
-}
+inline bool tensorCoresAvailable() { return DeviceInfoCache::instance().hasTensorCores(); }
 
 /**
  * 检查给定维度是否适合 Tensor Core 加速
@@ -67,11 +61,7 @@ inline bool tensorCoreDimensionsSupported(int M, int K, int N) {
  * 获取当前设备的 Tensor Core 信息字符串
  */
 inline const char *getTensorCoreArchName() {
-    int device;
-    CUDA_CHECK(cudaGetDevice(&device));
-
-    cudaDeviceProp prop;
-    CUDA_CHECK(cudaGetDeviceProperties(&prop, device));
+    const cudaDeviceProp &prop = DeviceInfoCache::instance().prop();
 
     if (prop.major == 7) {
         return (prop.minor == 0) ? "Volta" : (prop.minor == 5) ? "Turing" : "Unknown sm_7x";
@@ -101,7 +91,8 @@ using FallbackKernel =
  *
  * 提供一个空的 fallback（用于测试或显式配置场景）
  */
-inline void nullFallback(const float *, const float *, float *, int, int, int, cudaStream_t = 0) {
+[[maybe_unused]] inline void
+nullFallback(const float *, const float *, float *, int, int, int, cudaStream_t = 0) {
     // 空实现 - 用于测试
 }
 
@@ -179,7 +170,7 @@ __global__ void tensor_core_sgemm_kernel_fp16(const half *__restrict__ A,
  */
 inline void launch_tensor_core_sgemm_fp16_fast_path(const half *A, const half *B, float *C, int M,
                                                     int K, int N, cudaStream_t stream = 0) {
-    dim3 blockDim(32, 1);
+    dim3 blockDim(kDefaultTileSize, 1);
     dim3 gridDim((N + WMMA_N - 1) / WMMA_N, (M + WMMA_M - 1) / WMMA_M);
 
     tensor_core_sgemm_kernel_fp16<<<gridDim, blockDim, 0, stream>>>(A, B, C, M, K, N);
@@ -255,9 +246,17 @@ inline void launch_tensor_core_sgemm_with_fallback(const float *A, const float *
     DeviceMemory<half> d_A_fp16(num_A);
     DeviceMemory<half> d_B_fp16(num_B);
 
-    int blockSize = 256;
-    int gridSizeA = static_cast<int>((num_A + blockSize - 1) / blockSize);
-    int gridSizeB = static_cast<int>((num_B + blockSize - 1) / blockSize);
+    int blockSize = kDefaultBlockSize;
+    // 安全计算 gridSize，检查溢出
+    auto safeGridSize = [](size_t num, int blk) -> int {
+        size_t grid = (num + blk - 1) / blk;
+        if (grid > static_cast<size_t>(INT_MAX)) {
+            throw CudaError("Grid size overflow: matrix too large for kernel launch");
+        }
+        return static_cast<int>(grid);
+    };
+    int gridSizeA = safeGridSize(num_A, blockSize);
+    int gridSizeB = safeGridSize(num_B, blockSize);
 
     float_to_half_kernel<<<gridSizeA, blockSize, 0, stream>>>(A, d_A_fp16.get(),
                                                               static_cast<int>(num_A));

src/utils/benchmark_metrics.cuh

@@ -57,32 +57,11 @@ inline PerformanceMetrics calculateSgemmMetrics(int M, int K, int N, float time_
  * - 时钟频率
  */
 inline float getTheoreticalPeakGflops() {
-    int device;
-    CUDA_CHECK(cudaGetDevice(&device));
-
-    cudaDeviceProp prop;
-    CUDA_CHECK(cudaGetDeviceProperties(&prop, device));
-
-    // 每个 SM 的核心数（基于架构）
-    // 参考: https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#compute-capabilities
-    int coresPerSM;
-    if (prop.major == 7) {
-        coresPerSM = 64; // Volta (sm_70, sm_72), Turing (sm_75)
-    } else if (prop.major == 8) {
-        coresPerSM = (prop.minor == 0 || prop.minor == 6)
-                         ? 64
-                         : 128; // A100/sm_80, A10G/sm_86: 64, others: 128
-    } else if (prop.major == 9) {
-        coresPerSM = 128; // Hopper (sm_90)
-    } else {
-        coresPerSM = 64; // 默认回退
-    }
-
-    // 时钟频率 (kHz -> GHz)
-    float clockGHz = static_cast<float>(prop.clockRate) / 1e6f;
+    DeviceInfoCache &cache = DeviceInfoCache::instance();
+    const cudaDeviceProp &prop = cache.prop();
 
     // 峰值 GFLOPS = SMs * cores/SM * 2 (FMA) * clock (GHz) * 1000 (MHz factor)
-    float peakGflops = prop.multiProcessorCount * coresPerSM * 2 * clockGHz * 1000;
+    float peakGflops = prop.multiProcessorCount * cache.coresPerSM() * 2 * cache.clockGHz() * 1000;
 
     return peakGflops;
 }
@@ -96,11 +75,7 @@ inline float getTheoreticalPeakGflops() {
  * - 内存总线宽度
  */
 inline float getTheoreticalPeakBandwidth() {
-    int device;
-    CUDA_CHECK(cudaGetDevice(&device));
-
-    cudaDeviceProp prop;
-    CUDA_CHECK(cudaGetDeviceProperties(&prop, device));
+    const cudaDeviceProp &prop = DeviceInfoCache::instance().prop();
 
     // 内存时钟频率 (Hz -> MHz)
     float memoryClockMHz = static_cast<float>(prop.memoryClockRate) / 1000.0f;
@@ -139,8 +114,12 @@ inline float calculateEfficiency(float actual_gflops, float peak_gflops) {
 
 /**
  * 计算带宽利用率（相对于理论峰值的百分比）
+ *
+ * 注意：此函数为工具函数，供外部调用者使用。
+ * 内部 benchmark 流程使用 calculateEfficiency。
  */
-inline float calculateBandwidthUtilization(float actual_bandwidth, float peak_bandwidth) {
+[[maybe_unused]] inline float
+calculateBandwidthUtilization(float actual_bandwidth, float peak_bandwidth) {
     if (peak_bandwidth <= 0)
         return 0.0f;
     return (actual_bandwidth / peak_bandwidth) * 100.0f;
@@ -153,12 +132,16 @@ inline float calculateBandwidthUtilization(float actual_bandwidth, float peak_ba
 /**
  * 打印性能比较报告
  *
+ * 注意：此函数为工具函数，供外部调用者打印格式化报告。
+ * 内部 benchmark 流程使用 SGEMMBenchmark::printSummary。
+ *
  * @param kernel_name 内核名称
  * @param metrics 性能指标
  * @param baseline_gflops 基线 GFLOPS（如 cuBLAS）
  */
-inline void printPerformanceReport(const char *kernel_name, const PerformanceMetrics &metrics,
-                                   float baseline_gflops = 0.0f) {
+[[maybe_unused]] inline void
+printPerformanceReport(const char *kernel_name, const PerformanceMetrics &metrics,
+                       float baseline_gflops = 0.0f) {
     printf("  %-30s | %8.3f ms | %10.2f GFLOPS | %8.2f GB/s | AI: %.1f\n", kernel_name,
            metrics.time_ms, metrics.gflops, metrics.bandwidth_gb_s, metrics.arithmetic_intensity);
 

src/utils/cuda_utils.cuh

@@ -7,6 +7,25 @@
 #include <curand.h>
 #include <random>
 
+// ============================================================================
+// 命名常量
+// ============================================================================
+
+namespace config {
+/// 默认 tile 大小（用于 SGEMM 内核）
+inline constexpr int kDefaultTileSize = 32;
+
+/// 默认 block 大小（用于 CUDA 内核启动）
+inline constexpr int kDefaultBlockSize = 256;
+
+/// 文件名缓冲区大小
+inline constexpr int kFilenameBufferSize = 256;
+} // namespace config
+
+using config::kDefaultBlockSize;
+using config::kDefaultTileSize;
+using config::kFilenameBufferSize;
+
 // ============================================================================
 // Error Checking Macros
 // ============================================================================
@@ -56,7 +75,7 @@ template <typename T> class DeviceMemory {
   public:
     DeviceMemory() : ptr_(nullptr), size_(0) {}
 
-    explicit DeviceMemory(size_t count) : size_(count) {
+    explicit DeviceMemory(size_t count) : ptr_(nullptr), size_(count) {
         CUDA_CHECK(cudaMalloc(&ptr_, count * sizeof(T)));
     }
 
@@ -169,3 +188,62 @@ inline void printGPUInfo() {
     printf("  L2 Cache Size: %d KB\n", prop.l2CacheSize / 1024);
     printf("\n");
 }
+
+// ============================================================================
+// Device Info Cache - 缓存设备属性避免重复查询
+// ============================================================================
+
+/**
+ * 设备信息缓存类（单例模式）
+ *
+ * 缓存 cudaDeviceProp 和常用计算值，避免重复调用 cudaGetDeviceProperties。
+ * 首次访问时初始化，之后返回缓存值。
+ */
+class DeviceInfoCache {
+  public:
+    /// 获取单例实例
+    static DeviceInfoCache &instance() {
+        static DeviceInfoCache cache;
+        return cache;
+    }
+
+    /// 获取缓存的设备属性
+    const cudaDeviceProp &prop() const { return prop_; }
+
+    /// 获取设备 ID
+    int deviceId() const { return device_; }
+
+    /// 检查是否支持 Tensor Core (sm_70+)
+    bool hasTensorCores() const { return prop_.major >= 7; }
+
+    /// 获取每个 SM 的 CUDA 核心数（基于架构）
+    int coresPerSM() const { return coresPerSM_; }
+
+    /// 获取时钟频率 (GHz)
+    float clockGHz() const { return clockGHz_; }
+
+  private:
+    DeviceInfoCache() {
+        CUDA_CHECK(cudaGetDevice(&device_));
+        CUDA_CHECK(cudaGetDeviceProperties(&prop_, device_));
+        coresPerSM_ = computeCoresPerSM();
+        clockGHz_ = static_cast<float>(prop_.clockRate) / 1e6f;
+    }
+
+    int computeCoresPerSM() const {
+        // 参考: https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#compute-capabilities
+        if (prop_.major == 7) {
+            return 64; // Volta (sm_70, sm_72), Turing (sm_75)
+        } else if (prop_.major == 8) {
+            return (prop_.minor == 0 || prop_.minor == 6) ? 64 : 128; // A100/sm_80, A10G/sm_86: 64
+        } else if (prop_.major == 9) {
+            return 128; // Hopper (sm_90)
+        }
+        return 64; // 默认回退
+    }
+
+    int device_;
+    cudaDeviceProp prop_;
+    int coresPerSM_;
+    float clockGHz_;
+};