fix: update version, fix typos, and improve docs

- Fix CMakeLists.txt version: 2.0.0 → 2.1.0
- Fix Chinese typo in RELEASE_NOTES.md: "工你流" → "工作流"
- Update _config.yml: add SEO config, social links, better description
- Update index.md: add description, forks badge, improve performance table
- Update ci.yml: add sm_89, sm_90 architectures, improve info output
- Improve benchmark.cuh: add comments for GPU architecture cores
- Refactor test_sgemm.cu: use RAII DeviceMemory wrapper instead of raw pointers

Author: LessUp

.github/workflows/ci.yml

@@ -51,12 +51,13 @@ jobs:
           cmake -S . -B build
           -DCMAKE_BUILD_TYPE=Release
           -DBUILD_TESTS=OFF
-          -DCMAKE_CUDA_ARCHITECTURES="70;75;80;86"
+          -DCMAKE_CUDA_ARCHITECTURES="70;75;80;86;89;90"
 
       - name: Build
         run: cmake --build build --target sgemm_benchmark -j$(nproc)
 
       - name: Info
         run: |
           echo "✅ CUDA compilation successful"
+          echo "ℹ️  Supported architectures: sm_70, sm_75, sm_80, sm_86, sm_89, sm_90"
           echo "ℹ️  GPU runtime tests require a CUDA-capable machine"

CMakeLists.txt

@@ -1,6 +1,6 @@
 cmake_minimum_required(VERSION 3.18)
 project(sgemm_optimization
-    VERSION 2.0.0
+    VERSION 2.1.0
     DESCRIPTION "SGEMM optimization from naive to Tensor Core"
     LANGUAGES CXX CUDA
 )

RELEASE_NOTES.md

@@ -37,7 +37,7 @@
     - 🚀 快速开始指南
     - 📖 GPU 架构支持表
   - 优化 `_config.yml` Jekyll 配置
-  - **GitHub Workflows**: 简化并优化工你流
+  - **GitHub Workflows**: 简化并优化工作流
     - `ci.yml`: 清晰的步骤命名
     - `pages.yml`: 修复 paths 过滤，改进并发配置
 

_config.yml

@@ -10,24 +10,38 @@ plugins:
   - jekyll-remote-theme
   - jekyll-seo-tag
   - jekyll-sitemap
+  - jekyll-default-layout
 
 # Site info
 title: SGEMM Optimization
+tagline: From Naive to Tensor Core
 description: >-
-  Hand-written, progressively optimized CUDA matrix multiplication — the "Hello World" of HPC.
-  Five kernel variants from naive triple-loop to Tensor Core WMMA.
+  Progressive CUDA SGEMM optimization with 5 kernel variants demonstrating
+  GPU programming techniques including WMMA, shared memory tiling,
+  bank conflict elimination, and double buffering.
 repository: LessUp/sgemm-optimization
+author: LessUp
+
+# SEO and Social
+social:
+  name: LessUp
+  links:
+    - https://github.com/LessUp
 
 # URLs
 url: "https://lessup.github.io"
 baseurl: /sgemm-optimization
 
 # Localization
 lang: en
-author: LessUp
 
 # Navigation
 show_downloads: true
+github:
+  is_project_page: true
+  repository_url: https://github.com/LessUp/sgemm-optimization
+  zip_url: https://github.com/LessUp/sgemm-optimization/zipball/master
+  tar_url: https://github.com/LessUp/sgemm-optimization/tarball/master
 
 # Markdown
 markdown: kramdown

index.md

@@ -1,6 +1,8 @@
 ---
 layout: default
 title: SGEMM Optimization
+nav_exclude: false
+description: Progressive CUDA SGEMM optimization with 5 kernel variants demonstrating GPU programming techniques including WMMA, shared memory tiling, bank conflict elimination, and double buffering.
 ---
 
 # SGEMM Optimization: From Naive to Tensor Core
@@ -11,6 +13,7 @@ title: SGEMM Optimization
 ![CUDA](https://img.shields.io/badge/CUDA-11.0+-76B900?logo=nvidia&logoColor=white)
 ![C++](https://img.shields.io/badge/C%2B%2B-17-00599C?logo=c%2B%2B&logoColor=white)
 [![GitHub stars](https://img.shields.io/github/stars/LessUp/sgemm-optimization?style=social)](https://github.com/LessUp/sgemm-optimization/stargazers)
+[![GitHub forks](https://img.shields.io/github/forks/LessUp/sgemm-optimization?style=social)](https://github.com/LessUp/sgemm-optimization/network/members)
 
 **Hand-written, progressively optimized CUDA matrix multiplication — the "Hello World" of HPC**
 
@@ -39,18 +42,18 @@ cmake --build build -j$(nproc)
 
 ## 📊 Performance Summary
 
-NVIDIA RTX 3060 Laptop (Ampere, sm_86), 1024×1024×1024:
+Tested on NVIDIA RTX 3060 Laptop (Ampere, sm_86), Matrix Size: 1024×1024×1024
 
 | Kernel | GFLOPS | vs cuBLAS | Time | Key Technique |
 |--------|-------:|----------:|------:|---------------|
-| **cuBLAS** | 5727 | 100% | 0.375 ms | NVIDIA library |
-| **Tensor Core** | 2300 | 40.2% | 0.934 ms | WMMA API, FP16 |
-| **Tiled** | 753 | 13.1% | 2.853 ms | Shared memory |
-| **Double Buffer** | 701 | 12.2% | 3.064 ms | Pipeline overlap |
-| **Bank-Free** | 673 | 11.8% | 3.190 ms | Padding |
-| **Naive** | 604 | 10.6% | 3.553 ms | Baseline |
-
-> Performance varies by GPU, CUDA version, and matrix size. [See full benchmarks](docs/benchmark-results.md).
+| **cuBLAS** | 5727 | 100% | 0.375 ms | NVIDIA optimized library |
+| **Tensor Core** | 2300 | 40.2% | 0.934 ms | WMMA API, FP16 compute |
+| **Tiled** | 753 | 13.1% | 2.853 ms | Shared memory blocking |
+| **Double Buffer** | 701 | 12.2% | 3.064 ms | Latency hiding pipeline |
+| **Bank-Free** | 673 | 11.8% | 3.190 ms | Bank conflict elimination |
+| **Naive** | 604 | 10.6% | 3.553 ms | Baseline implementation |
+
+> ⚠️ Performance varies by GPU model, CUDA version, and matrix size. [See full benchmark analysis](docs/benchmark-results.md).
 
 ---
 

src/utils/benchmark.cuh

@@ -319,13 +319,13 @@ inline float getTheoreticalPeakGflops() {
   // See: https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#compute-capabilities
   int coresPerSM;
   if (prop.major == 7) {
-    coresPerSM = (prop.minor == 0 || prop.minor == 2) ? 64 : 64; // Volta, Turing
+    coresPerSM = 64; // Volta (sm_70, sm_72), Turing (sm_75)
   } else if (prop.major == 8) {
-    coresPerSM = (prop.minor == 0 || prop.minor == 6) ? 64 : 128; // Ampere
+    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
+    coresPerSM = 128; // Hopper (sm_90), Ada (sm_89 is major=8)
   } else {
-    coresPerSM = 64; // Default fallback
+    coresPerSM = 64; // Default fallback for older architectures
   }
 
   // Use actual GPU clock rate (convert from kHz to GHz)

tests/test_sgemm.cu

@@ -8,6 +8,7 @@
 
 #include <cuda_runtime.h>
 #include <gtest/gtest.h>
+#include <memory>
 #include <random>
 #include <tuple>
 #include <vector>
@@ -147,54 +148,42 @@ protected:
     initRandomMatrix(h_A_.data(), M_, K_, -1.0f, 1.0f, 42);
     initRandomMatrix(h_B_.data(), K_, N_, -1.0f, 1.0f, 123);
 
-    CUDA_CHECK(cudaMalloc(&d_A_, M_ * K_ * sizeof(float)));
-    CUDA_CHECK(cudaMalloc(&d_B_, K_ * N_ * sizeof(float)));
-    CUDA_CHECK(cudaMalloc(&d_C_, M_ * N_ * sizeof(float)));
-    CUDA_CHECK(cudaMalloc(&d_C_ref_, M_ * N_ * sizeof(float)));
+    d_A_ = std::make_unique<DeviceMemory<float>>(M_ * K_);
+    d_B_ = std::make_unique<DeviceMemory<float>>(K_ * N_);
+    d_C_ = std::make_unique<DeviceMemory<float>>(M_ * N_);
+    d_C_ref_ = std::make_unique<DeviceMemory<float>>(M_ * N_);
 
-    CUDA_CHECK(cudaMemcpy(d_A_, h_A_.data(), M_ * K_ * sizeof(float), cudaMemcpyHostToDevice));
-    CUDA_CHECK(cudaMemcpy(d_B_, h_B_.data(), K_ * N_ * sizeof(float), cudaMemcpyHostToDevice));
+    d_A_->copyFromHost(h_A_.data(), M_ * K_);
+    d_B_->copyFromHost(h_B_.data(), K_ * N_);
 
-    verifier_.computeReference(d_A_, d_B_, d_C_ref_, M_, K_, N_);
-    CUDA_CHECK(
-        cudaMemcpy(h_C_ref_.data(), d_C_ref_, M_ * N_ * sizeof(float), cudaMemcpyDeviceToHost));
-  }
-
-  void TearDown() override {
-    if (d_A_)
-      cudaFree(d_A_);
-    if (d_B_)
-      cudaFree(d_B_);
-    if (d_C_)
-      cudaFree(d_C_);
-    if (d_C_ref_)
-      cudaFree(d_C_ref_);
+    verifier_.computeReference(d_A_->get(), d_B_->get(), d_C_ref_->get(), M_, K_, N_);
+    d_C_ref_->copyToHost(h_C_ref_.data(), M_ * N_);
   }
 
   template <typename LaunchFn>
   VerifyResult runKernelAndCompare(LaunchFn launch_fn,
                                    VerifyTolerance tolerance = kStandardVerifyTolerance) {
-    CUDA_CHECK(cudaMemset(d_C_, 0, M_ * N_ * sizeof(float)));
+    d_C_->zero();
     launch_fn();
     CUDA_CHECK(cudaDeviceSynchronize());
-    CUDA_CHECK(cudaMemcpy(h_C_.data(), d_C_, M_ * N_ * sizeof(float), cudaMemcpyDeviceToHost));
+    d_C_->copyToHost(h_C_.data(), M_ * N_);
     return compareMatrices(h_C_.data(), h_C_ref_.data(), M_, N_, tolerance);
   }
 
   int M_, K_, N_;
   std::vector<float> h_A_, h_B_, h_C_, h_C_ref_;
-  float *d_A_ = nullptr;
-  float *d_B_ = nullptr;
-  float *d_C_ = nullptr;
-  float *d_C_ref_ = nullptr;
+  std::unique_ptr<DeviceMemory<float>> d_A_;
+  std::unique_ptr<DeviceMemory<float>> d_B_;
+  std::unique_ptr<DeviceMemory<float>> d_C_;
+  std::unique_ptr<DeviceMemory<float>> d_C_ref_;
   SGEMMVerifier verifier_;
 };
 
 class NaiveSGEMMTest : public SGEMMKernelTest {};
 
 TEST_P(NaiveSGEMMTest, CorrectnessProperty) {
   VerifyResult result =
-      runKernelAndCompare([&] { launch_naive_sgemm<>(d_A_, d_B_, d_C_, M_, K_, N_); });
+      runKernelAndCompare([&] { launch_naive_sgemm<>(d_A_->get(), d_B_->get(), d_C_->get(), M_, K_, N_); });
 
   EXPECT_TRUE(result.passed) << "Naive SGEMM failed for dimensions " << M_ << "x" << K_ << "x" << N_
                              << " (max_rel_error: " << result.max_rel_error << ")";
@@ -207,7 +196,7 @@ class TiledSGEMMTest : public SGEMMKernelTest {};
 
 TEST_P(TiledSGEMMTest, CorrectnessProperty) {
   VerifyResult result =
-      runKernelAndCompare([&] { launch_tiled_sgemm<32>(d_A_, d_B_, d_C_, M_, K_, N_); });
+      runKernelAndCompare([&] { launch_tiled_sgemm<32>(d_A_->get(), d_B_->get(), d_C_->get(), M_, K_, N_); });
 
   EXPECT_TRUE(result.passed) << "Tiled SGEMM failed for dimensions " << M_ << "x" << K_ << "x" << N_
                              << " (max_rel_error: " << result.max_rel_error << ")";
@@ -220,7 +209,7 @@ class BankConflictFreeSGEMMTest : public SGEMMKernelTest {};
 
 TEST_P(BankConflictFreeSGEMMTest, CorrectnessProperty) {
   VerifyResult result = runKernelAndCompare(
-      [&] { launch_bank_conflict_free_sgemm<32>(d_A_, d_B_, d_C_, M_, K_, N_); });
+      [&] { launch_bank_conflict_free_sgemm<32>(d_A_->get(), d_B_->get(), d_C_->get(), M_, K_, N_); });
 
   EXPECT_TRUE(result.passed) << "BankConflictFree SGEMM failed for dimensions " << M_ << "x" << K_
                              << "x" << N_ << " (max_rel_error: " << result.max_rel_error << ")";
@@ -233,7 +222,7 @@ class DoubleBufferSGEMMTest : public SGEMMKernelTest {};
 
 TEST_P(DoubleBufferSGEMMTest, CorrectnessProperty) {
   VerifyResult result =
-      runKernelAndCompare([&] { launch_double_buffer_sgemm<32>(d_A_, d_B_, d_C_, M_, K_, N_); });
+      runKernelAndCompare([&] { launch_double_buffer_sgemm<32>(d_A_->get(), d_B_->get(), d_C_->get(), M_, K_, N_); });
 
   EXPECT_TRUE(result.passed) << "DoubleBuffer SGEMM failed for dimensions " << M_ << "x" << K_
                              << "x" << N_ << " (max_rel_error: " << result.max_rel_error << ")";
@@ -252,7 +241,7 @@ TEST_P(TensorCoreSGEMMTest, FastPathCorrectnessProperty) {
   ASSERT_TRUE(tensorCoreDimensionsSupported(M_, K_, N_));
 
   VerifyResult result = runKernelAndCompare(
-      [&] { launch_tensor_core_sgemm(d_A_, d_B_, d_C_, M_, K_, N_); }, kTensorCoreVerifyTolerance);
+      [&] { launch_tensor_core_sgemm(d_A_->get(), d_B_->get(), d_C_->get(), M_, K_, N_); }, kTensorCoreVerifyTolerance);
 
   EXPECT_TRUE(result.passed) << "TensorCore SGEMM fast path failed for dimensions " << M_ << "x"
                              << K_ << "x" << N_ << " (max_rel_error: " << result.max_rel_error
@@ -266,7 +255,7 @@ class TensorCoreFallbackTest : public SGEMMKernelTest {};
 
 TEST_P(TensorCoreFallbackTest, NonAlignedInputsFallbackSafely) {
   VerifyResult result = runKernelAndCompare(
-      [&] { launch_tensor_core_sgemm(d_A_, d_B_, d_C_, M_, K_, N_); }, kStandardVerifyTolerance);
+      [&] { launch_tensor_core_sgemm(d_A_->get(), d_B_->get(), d_C_->get(), M_, K_, N_); }, kStandardVerifyTolerance);
 
   EXPECT_TRUE(result.passed) << "TensorCore fallback failed for dimensions " << M_ << "x" << K_
                              << "x" << N_ << " (max_rel_error: " << result.max_rel_error << ")";
@@ -303,26 +292,22 @@ TEST_F(DimensionInvarianceTest, AllStandardKernelsWorkWithVariousDimensions) {
     initRandomMatrix(h_A.data(), M, K, -1.0f, 1.0f, iter);
     initRandomMatrix(h_B.data(), K, N, -1.0f, 1.0f, iter + 1000);
 
-    float *d_A = nullptr;
-    float *d_B = nullptr;
-    float *d_C = nullptr;
-    float *d_ref = nullptr;
-    CUDA_CHECK(cudaMalloc(&d_A, M * K * sizeof(float)));
-    CUDA_CHECK(cudaMalloc(&d_B, K * N * sizeof(float)));
-    CUDA_CHECK(cudaMalloc(&d_C, M * N * sizeof(float)));
-    CUDA_CHECK(cudaMalloc(&d_ref, M * N * sizeof(float)));
+    DeviceMemory<float> d_A(M * K);
+    DeviceMemory<float> d_B(K * N);
+    DeviceMemory<float> d_C(M * N);
+    DeviceMemory<float> d_ref(M * N);
 
-    CUDA_CHECK(cudaMemcpy(d_A, h_A.data(), M * K * sizeof(float), cudaMemcpyHostToDevice));
-    CUDA_CHECK(cudaMemcpy(d_B, h_B.data(), K * N * sizeof(float), cudaMemcpyHostToDevice));
+    d_A.copyFromHost(h_A.data(), M * K);
+    d_B.copyFromHost(h_B.data(), K * N);
 
-    computeReference(handle_, d_A, d_B, d_ref, M, K, N);
-    CUDA_CHECK(cudaMemcpy(h_ref.data(), d_ref, M * N * sizeof(float), cudaMemcpyDeviceToHost));
+    computeReference(handle_, d_A.get(), d_B.get(), d_ref.get(), M, K, N);
+    d_ref.copyToHost(h_ref.data(), M * N);
 
     auto testKernel = [&](const char *name, auto kernel_func) {
-      CUDA_CHECK(cudaMemset(d_C, 0, M * N * sizeof(float)));
-      kernel_func(d_A, d_B, d_C, M, K, N);
+      d_C.zero();
+      kernel_func(d_A.get(), d_B.get(), d_C.get(), M, K, N);
       CUDA_CHECK(cudaDeviceSynchronize());
-      CUDA_CHECK(cudaMemcpy(h_C.data(), d_C, M * N * sizeof(float), cudaMemcpyDeviceToHost));
+      d_C.copyToHost(h_C.data(), M * N);
 
       VerifyResult result =
           compareMatrices(h_C.data(), h_ref.data(), M, N, kStandardVerifyTolerance);
@@ -334,11 +319,6 @@ TEST_F(DimensionInvarianceTest, AllStandardKernelsWorkWithVariousDimensions) {
     testKernel("Tiled", launch_tiled_sgemm<32>);
     testKernel("BankConflictFree", launch_bank_conflict_free_sgemm<32>);
     testKernel("DoubleBuffer", launch_double_buffer_sgemm<32>);
-
-    cudaFree(d_A);
-    cudaFree(d_B);
-    cudaFree(d_C);
-    cudaFree(d_ref);
   }
 }